Threads, blocks, and warps

A CUDA program runs on a heterogeneous system: a host (the CPU) and a device (the GPU), connected by an interconnect such as PCIe or NVLink. The host code copies data to device memory, launches a kernel, and waits for it to complete. A kernel launch starts many threads, often millions, all executing the same device code.

Those threads are organized into thread blocks, and thread blocks into a grid. Every thread block in a grid runs entirely on one streaming multiprocessor (SM), which is what lets threads inside a block synchronize and share on-chip memory cheaply. There is no such guarantee across blocks: the CUDA programming model requires that thread blocks be safe to run in any order, in parallel or in series, because a grid can have far more blocks than the GPU has SMs to run them on at once.

Inside a block, threads execute in fixed groups of 32 called warps, in a Single-Instruction Multiple-Threads (SIMT) model: every thread in a warp runs the same instruction at the same time, but each thread carries its own program counter and can take a different branch. When threads in a warp disagree on which branch to take, the ones not on the active path are masked off until the warp reconverges, a cost called warp divergence. It follows that a block sized to a multiple of 32 threads uses every lane of its last warp; anything else leaves lanes idle for the whole kernel.

In code, the whole model fits in a dozen lines. A kernel is a __global__ function; the triple-chevron launch names the grid and block dimensions; and inside the kernel, each thread combines threadIdx, blockIdx, and blockDim to find the one element it is responsible for. This is the CUDA programming guide’s own first example, an element-wise vector addition where every thread performs exactly one add:

__global__ void vecAdd(float* A, float* B, float* C)
{
    int workIndex = threadIdx.x + blockDim.x * blockIdx.x;
    C[workIndex] = A[workIndex] + B[workIndex];
}

int main()
{
    // ...
    vecAdd<<<1, 256>>>(A, B, C);
    // ...
}

The launch <<<1, 256>>> starts one thread block of 256 threads, and the guide notes the two constraints this book keeps returning to: a block may contain at most 1,024 threads because the whole block must fit on one SM, and kernel launches are asynchronous, so the host must synchronize before it can trust the result.

One block of 256 threads covers 256 elements, but the same index expression scales to any number of blocks. Launched as vecAdd<<<4, 256>>> over a vector of 1,024 elements, blockDim.x * blockIdx.x becomes each block’s offset into the vector: threads in the first block compute indices 0 through 255, threads in the second land at threadIdx.x + 256, the third at threadIdx.x + 512. Real vector lengths are not always multiples of the block size, so the guide’s full kernel takes the length as a parameter and guards the work with if (workIndex < vectorLength); threads past the end simply do nothing. The launch then rounds the block count up with an integer ceiling divide, (vectorLength + threads - 1) / threads. A few idle threads in the last block cost little, the guide notes, but launching whole blocks in which no thread does work should be avoided.

None of this runs until the arrays live in memory the GPU can reach. The explicit path allocates device buffers with cudaMalloc and copies data across with cudaMemcpy, whose last argument names the direction: cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, or cudaMemcpyDefault, which infers the direction from the pointer values. Wrapped around the launch, those calls turn vecAdd into a complete program:

cudaMalloc(&devA, vectorLength*sizeof(float));
cudaMalloc(&devB, vectorLength*sizeof(float));
cudaMalloc(&devC, vectorLength*sizeof(float));

cudaMemcpy(devA, A, vectorLength*sizeof(float), cudaMemcpyDefault);
cudaMemcpy(devB, B, vectorLength*sizeof(float), cudaMemcpyDefault);

int threads = 256;
int blocks = cuda::ceil_div(vectorLength, threads);
vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);
// wait for kernel execution to complete
cudaDeviceSynchronize();

// Copy results back to host
cudaMemcpy(C, devC, vectorLength*sizeof(float), cudaMemcpyDefault);

cudaFree(devA);
cudaFree(devB);
cudaFree(devC);

Two details in that listing carry most of the meaning. cudaMemcpy is synchronous: it does not return until the copy has completed. The kernel launch is not, which is why cudaDeviceSynchronize sits between the launch and the copy back: it blocks the host thread until all previously issued GPU work has finished. The guide also offers a second path, unified memory, where cudaMallocManaged allocates buffers the driver keeps accessible to both CPU and GPU and the copies disappear from the source. The explicit version is more verbose precisely because it affords control over when data moves and where it lives, control the performance chapters of this book spend heavily.