AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter
Files
cd091261309f96f2ea257490e2434fe73e86dc4e
rocm-systems/projects/hip/docs/markdown/modern_cpp_features.md
T
Finlay 35058813fb added document of modern cpp features supported by the compiler (#2848) 
[ROCm/hip commit: 02ee43c5b4]
2022-09-27 10:46:35 +05:30

9.8 KiB
Rå Blame Historik

HIP C++ Feature

C++ 11

Rvalue References

struct Y {
  int x;
};

__device__ void do_something(Y &&val) { val.x += 1; }

__global__ void kernel() {
  Y y{10};
  // do_something(y); // does not compile since the argument is an lvalue
  do_something(std::move(y));
}

int main() { kernel<<<1, 1>>>(); }

Rvalue References for *this

struct Sample {
  __host__ __device__ void callMe() & { printf("Lval Func\n"); }
  __host__ __device__ void callMe() && { printf("Rval Func\n"); }
};

__global__ void kernel() {
  Sample s;
  s.callMe();        // prints Lval Func
  Sample().callMe(); // prints Rval Func
}

int main() { kernel<<<1, 1>>>(); }

Variadic templates, Static Assertions, auto Variables

template <typename T> __host__ __device__ T add(T val) { return val; }

template <typename T, typename... Targs>
__host__ __device__ T add(T val, Targs... pVal) {
  static_assert(std::is_arithmetic<T>::value, "Not a valid type");
  return val + add(pVal...);
}

template <typename T, typename... Targs>
__global__ void kernel(T *ptr, Targs... args) {
  auto &&sum = add(args...);
  *ptr = sum;
}

// Or something like

__device__ int &getX(int &x) { return ++x; }
__device__ int getY(int &x) { return x + 10; }

__global__ void kernel() {
  int X = 0;
  auto &&x = getX(X);
  auto &&y = getY(X);

  // Init with value or initializer list
  auto val{10};
  auto list = {10};
}

int main() { kernel<<<1, 1>>>(); }

Non-static Data Member Initialization

struct S {
  int a = 1;
  int b = 2;
};

__global__ void kernel() {
  S s; // s.a == 1 and s.b == 2
}
int main() { kernel<<<1, 1>>>(); }

Lambda Device Functions

template <typename T> __global__ void kernel(T f) { f(); }

int main() {
  auto func = [=] __device__() { printf("In Kernel\n"); };
  kernel<<<1, 1>>>(func);
  hipDeviceSynchronize();
}

decltype Usage

template <typename T> __device__ T ret() {
  T x{0};
  return x;
}

template <typename T> __global__ void kernel() {
  decltype(ret<T>()) a;
  int i = 0;
  decltype(i) j = i + 1;
}
int main() { kernel<float><<<1, 1>>>(); }

Default Template Arguments

template <int N = 5> __global__ void kernel(int x) { x += N; }
int main() {
  kernel<<<1, 1>>>(1);
  kernel<-2><<<1, 1>>>(1);
}

Template Alias

template <typename T> struct Alloc {};

template <typename T, typename U> struct Vector {};

template <typename T> using V = Vector<T, Alloc<T>>;

template <typename T> __global__ void kernel(T x) { V<T> v; }

int main() { kernel<<<1, 1>>>(5); }

Extern Template

template <typename T> __global__ void kernel(T x) {}

extern template __global__ void kernel(long x);

int main() {
  kernel<<<1, 1, 0, 0>>>(10); // will create a template specialization
  // kernel<<<1,1,0,0>>>(10l);  // looks for existing kernel<long>, causing
  // linking to fail
}

nullptr as a Keyword in Device Compiler

__global__ void kernel() {
  int *ptr = nullptr;
  //...
}
int main() { kernel<<<1, 1>>>(); }

Strongly Typed Enums

enum class EnumVals { Red, Blue, Green };
__global__ void kernel() {
  auto val = EnumVals::Red;
  //...
}
int main() { kernel<<<1, 1>>>(); }

Standardized Attribute Syntax

[[deprecated]] __global__ void kernel() {
  //...
}
int main() { kernel<<<1, 1>>>(); }

constexpr

struct S {
  constexpr __device__ S(double v) : val(v) {}
  constexpr __device__ double value() const { return val; }

private:
  double val;
};

constexpr __device__ int factorial(int n) {
  return n <= 1 ? 1 : (n * factorial(n - 1));
}

__global__ void kernel() {
  constexpr S s(factorial(5));
  constexpr double d = s.value();
  // ...
}
int main() { kernel<<<1, 1>>>(); }

alignas with Struct

struct alignas(alignof(int)) S {
  //...
};

__global__ void kernel() {
  S s;
  static_assert(alignof(S) == alignof(int), "they have the same alignment");
  // check the alignment
}
int main() { kernel<<<1, 1>>>(); }

Delegating Constructors

struct S {
private:
  int val;

public:
  __device__ S(int v) : val(v) {}
  __device__ S() : S(42) {}
};
__global__ void kernel() { S s{}; }
int main() { kernel<<<1, 1>>>(); }

Explicit Conversion Functions

struct S {
private:
  int val;

public:
  __device__ S(int val) : val(val) {}
  __device__ explicit operator int *() { return &val; }
};

__global__ void kernel() {
  S s{0};
  // if (s) { // compile error
  // without the explicit function specifier then s would be converted to the
  // pointer to s.val, which would be non-zero so always true.
  //}
  if ((int *)(s)) {
    // this compiles but is likely not what the user intended
  }
}
int main() { kernel<<<1, 1>>>(); }

Unicode Character Types, Unicode String, Universal Character Literal

__global__ void kernel() {
  // cant print it since printf(gpu) doesnot support unicode char arguments
  char16_t a = u'y';
  char32_t l = U'猫';
  auto *string = U"इस अनुवाद को करने से आपको क्या मिला?";
}
int main() { kernel<<<1, 1>>>(); }

User Defined Literals

__device__ long double operator"" _w(long double a) { return a; }
__device__ unsigned operator"" _w(char const *c) { return *c - '0'; }

__global__ void kernel() {
  auto ld = 1.2_w; // calls operator "" _w(1.2L)
  auto val = 2_w;  // calls operator "" _w("2")
}
int main() { kernel<<<1, 1>>>(); }

default/delete Functions

struct S {
  __device__ S() = default;
  __device__ S &operator=(const S &) = delete;
};
__global__ void kernel() {
  S s, other; // fine
  // other = s; // compile error, function deleted
}
int main() { kernel<<<1, 1>>>(); }

Friend Declaration

struct Y {};

struct A {
  __device__ A() = default;
  friend Y;
  // friend Z; // compile error since class or struct Z doesn't exist
  friend class Z;        // this is fine
  friend void asdf(int); // functions can be declared without a definition
};
__global__ void kernel() { A a; }
int main() { kernel<<<1, 1>>>(); }

Extended sizeof

template <typename... Ts> __global__ void kernel(Ts... ts) {
  auto size = sizeof...(ts);
  // ...
}
int main() { kernel<<<1, 1>>>(); }

Unrestricted Unions

struct Point {
  __device__ Point() {}
  __device__ Point(int x, int y) : x_(x), y_(y) {}
  int x_, y_;
};

union U {
  int z;
  double w;
  Point p;
  __device__ U() {}
  __device__ U(const Point &pt) : p(pt) {}
  __device__ U &operator=(const Point &pt) {
    new (&p) Point(pt);
    return *this;
  }
};

__global__ void kernel() {
  U u;
  //...
}
int main() { kernel<<<1, 1>>>(); }

Inline Namespaces

namespace XX {
inline namespace YY {
struct Y {
  int x;
};
} // namespace YY
struct X {
  int a;
};
} // namespace XX

__global__ void kernel() {
  XX::X x{};
  XX::Y y{};
}

int main() { kernel<<<1, 1>>>(); }

Range Based For-loop

__global__ void kernel() {
  for (auto &x : {1, 2, 3, 4, 5}) {
    // ...
  }
}
int main() { kernel<<<1, 1>>>(); }

override Specifier

struct Base {
  int n;
  __device__ Base(int v) : n(v + 1) {}
  __device__ Base() : Base(10) {}
  __device__ virtual ~Base() {}
  __device__ virtual int get() { return n; }
};

struct Derived : public Base {
  int n;
  __device__ Derived(int v) : n(v) {}
  __device__ int get() override { return n; }
  __device__ ~Derived() {}
};

__global__ void kernel() {
  Derived d(10);
  //...
}
int main() { kernel<<<1, 1>>>(); }

noexcept Keyword

__global__ void kernel() noexcept {
  int n;
  //...
}
int main() { kernel<<<1, 1>>>(); }

Consecutive Right Angle Brackets in Templates

template <typename T> struct A { T a; };

template <typename T> struct B { T b; };

__global__ void kernel() {
  A<B<int>> ab;
  //...
}
int main() { kernel<<<1, 1>>>(); }

Not Yet Documented

C++14

Not Yet Documented

Referens i nytt ärende View Git Blame Copy Permalink