CUDA C++

C++17/20 in device code (with limits): templates, constexpr, auto, lambdas, classes with __device__ methods, operator overloading, if constexpr, variadic templates.

29.1 Templates & generic kernels

template <typename T, int BLOCK>
__global__ void scaleKernel(T* data, T factor, int n) {
    int i = blockIdx.x * BLOCK + threadIdx.x;
    if (i < n) data[i] *= factor;
}
scaleKernel<float, 256><<<(n + 255) / 256, 256>>>(d_f, 2.0f, n);
scaleKernel<double, 128><<<(n + 127) / 128, 128>>>(d_d, 2.0, n);

29.2 Functors & lambdas

template <typename F>
__global__ void transformKernel(float* x, int n, F f) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) x[i] = f(x[i]);
}
auto square = [] __device__ (float v) { return v * v; };    // needs --extended-lambda
transformKernel<<<blocks, threads>>>(d_x, n, square);

29.3 Supported C++ features

C++17/20 in device code (with limits): templates, constexpr, auto, lambdas, classes with __device__ methods, operator overloading, if constexpr, variadic templates. Not supported in device code: exceptions, RTTI (dynamic_cast, typeid), standard library containers (std::vector) — use libcu++, Thrust, or custom types. Virtual functions work only for objects created on the device.

29.4 libcu++ (CUDA C++ Standard Library) & CCCL

CCCL (CUDA Core Compute Libraries) unifies Thrust, CUB, and libcu++, shipped with the toolkit and on GitHub (NVIDIA/cccl). cuda::std::mdspan, cuda::memory_resource, and C++ parallel algorithms with nvc++ -stdpar (§37) are part of the modern CUDA C++ stack.

#include <cuda/std/array>
#include <cuda/std/atomic>
#include <cuda/std/complex>
#include <cuda/std/span>
#include <cuda/atomic>
#include <cuda/barrier>

__device__ cuda::atomic<int, cuda::thread_scope_device> counter;
cuda::std::array<float, 4> v{1, 2, 3, 4};