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(std::move(y));
}
int main() { kernel<<<1, 1>>>(); }
Rvalue References for <tt>*this</tt>
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();
Sample().callMe();
}
int main() { kernel<<<1, 1>>>(); }
Variadic templates, Static Assertions, <tt>auto</tt> 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;
}
__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);
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;
}
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);
}
hipError_t hipDeviceSynchronize(void)
Waits on all active streams on current device.
<tt>decltype</tt> 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);
}
<tt>nullptr</tt> 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>>>(); }
<tt>constexpr</tt>
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>>>(); }
<tt>alignas</tt> with Struct
struct alignas(alignof(int)) S {
};
__global__ void kernel() {
S s;
static_assert(alignof(S) == alignof(int), "they have the same 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 ((int *)(s)) {
}
}
int main() { kernel<<<1, 1>>>(); }
Unicode Character Types, Unicode String, Universal Character Literal
__global__ void kernel() {
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;
auto val = 2_w;
}
int main() { kernel<<<1, 1>>>(); }
<tt>default</tt>/<tt>delete</tt> Functions
struct S {
__device__ S() = default;
__device__ S &operator=(const S &) = delete;
};
__global__ void kernel() {
S s, other;
}
int main() { kernel<<<1, 1>>>(); }
Friend Declaration
struct Y {};
struct A {
__device__ A() = default;
friend Y;
friend class Z;
friend void asdf(int);
};
__global__ void kernel() { A a; }
int main() { kernel<<<1, 1>>>(); }
Extended <tt>sizeof</tt>
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;
};
}
struct X {
int a;
};
}
__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>>>(); }
<tt>override</tt> 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>>>(); }
<tt>noexcept</tt> 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