|
HIP: Heterogenous-computing Interface for Portability
|
|
HIP: Heterogenous-computing Interface for Portability
|
```cpp 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>>>(); } ```
*this```cpp 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>>>(); } ```
auto Variables```cpp 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>>>(); } ```
```cpp 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>>>(); } ```
```cpp 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```cpp 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>>>(); } ```
```cpp template <int n="5"> global void kernel(int x) { x += N; } int main() { kernel<<<1, 1>>>(1); kernel<-2><<<1, 1>>>(1); } ```
```cpp 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); } ```
```cpp 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```cpp global void kernel() { int *ptr = nullptr; //... } int main() { kernel<<<1, 1>>>(); } ```
```cpp enum class EnumVals { Red, Blue, Green }; global void kernel() { auto val = EnumVals::Red; //... } int main() { kernel<<<1, 1>>>(); } ```
```cpp [[deprecated]] global void kernel() { //... } int main() { kernel<<<1, 1>>>(); } ```
constexpr```cpp 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```cpp 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>>>(); } ```
```cpp 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>>>(); } ```
```cpp struct S { private: int val;
public: device S(int val) : val(val) {} device explicit operator int *() { return } };
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>>>(); } ```
```cpp 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>>>(); } ```
```cpp 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```cpp 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>>>(); } ```
```cpp 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>>>(); } ```
sizeof```cpp template <typename... Ts> global void kernel(Ts... ts) { auto size = sizeof...(ts); // ... } int main() { kernel<<<1, 1>>>(); } ```
```cpp 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>>>(); } ```
```cpp 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>>>(); } ```
```cpp global void kernel() { for (auto &x : {1, 2, 3, 4, 5}) { // ... } } int main() { kernel<<<1, 1>>>(); } ```
override Specifier```cpp 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```cpp global void kernel() noexcept { int n; //... } int main() { kernel<<<1, 1>>>(); } ```
```cpp 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>>>(); } ```
1.8.5