What CUDA C++ Really Is

It's C++ with three syntax extensions — and one philosophy

At first glance CUDA C++ looks like plain C/C++. You include headers, you write functions, you compile. But three extensions change the game completely:

Function qualifiers — __global__, __device__, __host__ — decide where a function runs and who may call it. A __global__ kernel runs on the GPU but is launched from CPU host code.
The triple-angle launch — kernel<<<grid, block>>>(args); — is compile-time metadata that the runtime converts into a command buffer for the GPU scheduler.
SIMT built-ins — threadIdx.x, blockIdx.x, blockDim.x, gridDim.x — give every thread its own coordinates so it knows which slice of work to do.

Otherwise you inherit the full C++17 toolchain: lambdas, templates, <algorithm>, std::vector on the host side. The twist is that device code is funneled through NVCC, which splits your translation unit, compiles GPU sections with PTX/SASS back-ends, then relinks everything into a single binary.

Why a separate compiler? Because GPUs aren't another CPU. They have their own ISA (PTX is the portable intermediate, SASS is the per-architecture binary), their own register file shape, their own memory hierarchy. NVCC is the bilingual translator that lets you write one source file but ship two compiled programs that meet at link time.

The philosophy: SIMT (Single Instruction, Multiple Threads). Thousands of threads run the same kernel body in lockstep, branching on their indices. You don't write a loop over elements; you write what one thread does, then launch enough threads to cover the data.

Code

Three function qualifiers — same syntax, different homes·cuda

// __host__ — runs on CPU, callable from CPU. Default if no qualifier.
__host__ float cpu_norm(const float* a, int n) {
    float s = 0.f;
    for (int i = 0; i < n; ++i) s += a[i] * a[i];
    return std::sqrt(s);
}

// __device__ — runs on GPU, callable only from GPU code.
__device__ float square(float x) { return x * x; }

// __global__ — runs on GPU, callable only from CPU. The kernel.
__global__ void norm_kernel(const float* a, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = square(a[i]);  // calls __device__ helper
}

Triple-angle launch syntax — what NVCC sees·cuda

int n = 1 << 20;       // 1,048,576 elements
int threads = 256;     // threads per block (must be multiple of 32)
int blocks  = (n + threads - 1) / threads;  // ceil division

norm_kernel<<<blocks, threads>>>(d_a, d_out, n);
cudaDeviceSynchronize();   // wait for GPU to finish

Exercise

Open NVIDIA's PTX ISA reference page in your browser and skim the table of contents. Don't try to learn it — just notice that PTX has its own load/store, branch, and arithmetic instructions, distinct from x86 or ARM. This is the assembly your CUDA C++ becomes. Bookmark the page; you'll come back when a kernel underperforms and you need to read the generated PTX.

It's C++ with three syntax extensions — and one philosophy

Code

External links

Exercise

Progress

Comments 0