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: all threads in a warp execute the same kernel code, but they need not follow the same execution path through it. 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.
Where the data lives#
A memory hierarchy is laid over that execution hierarchy, one scope at a time. The programming guide tabulates five memory types by scope and lifetime: registers and local memory belong to a single thread and last for the kernel; shared memory belongs to a thread block and lasts for the kernel; global and constant memory have grid scope and last as long as the application. Registers and shared memory are physically located on the SM; the other three live in device memory. The closer a space sits to the thread, the narrower its scope and the shorter its life.
Registers are the fastest and the least visible. They sit on the SM, hold thread-local values, and are allocated by the compiler rather than by the programmer. Accessing one generally consumes no extra clock cycles per instruction, though delays can still arise from read-after-write dependencies and from register bank conflicts, which an application has no direct control over. What makes registers a design constraint rather than a free resource is that the file is finite and shared: to schedule a block onto an SM, the registers each thread needs multiplied by the number of threads in the block must fit in the SM’s register file, and a block that asks for more than the file holds makes the kernel unlaunchable until the block gets smaller.
When the compiler runs out of registers, the overflow goes to local memory, whose name misleads. Local describes its logical scope, one thread, not its physical location: local memory resides in device memory, so its accesses carry the same latency and bandwidth as global memory accesses and answer to the same coalescing rules. The compiler places variables there when it cannot determine that an array is indexed with constant quantities, when a structure or array would consume too much register space, and whenever a kernel uses more registers than are available. That last case is register spilling: values that were on-chip get written out to global memory and read back later to make room for others. The best practices guide is blunt about the consequence, that the word local in the name does not imply faster access.
Shared memory is the one on-chip space a programmer places data in deliberately. It is visible to every thread in a block, physically located on each SM, and it persists for the duration of the kernel: a user-managed scratchpad rather than a cache. Because it is on-chip, its bandwidth is higher and its latency lower than local or global memory, provided threads do not collide on the same bank. It shares physical storage with the L1 cache inside the SM’s unified data cache, so a kernel that uses shared memory reduces the L1 available to it, and a kernel that uses none leaves the whole space to L1. Allocation is per block, not per thread.
Because the threads of a block write into that array and then read what their neighbors wrote, they need a barrier in between. __syncthreads() is that barrier: it blocks every thread in the block until all of them have reached the call.
// assuming blockDim.x is 128
__global__ void example_syncthreads(int* input_data, int* output_data)
{
__shared__ int shared_data[128];
shared_data[threadIdx.x] = input_data[blockDim.x*blockIdx.x + threadIdx.x];
// All threads synchronize, guaranteeing all writes to 'shared_data' are ordered
// before any thread is unblocked from '__syncthreads()':
__syncthreads();
// A single thread safely reads 'shared_data':
if (threadIdx.x == 0) {
float sum = 0;
for (int i = 0; i < blockDim.x; ++i) {
sum += shared_data[i];
}
output_data[blockIdx.x] = sum;
}
}One __shared__ array of 128 integers is allocated for the whole block, not one per thread, which is what makes the last step legal: after the barrier, a single thread reads all 128 values that 128 different threads wrote. Sizing it at compile time like this is the static form. A kernel can instead request shared memory at launch, as a byte count in the third argument of the execution configuration, and declare the array extern __shared__ with empty brackets.
Constant memory is the odd member of the set. It has grid scope and application lifetime, it is read-only from the kernel, and it is small: a device has 64 KB of it in total. It is cached, so a read costs a device memory read only on a miss, but its cost model depends on agreement inside the warp. Accesses to different addresses by threads within a warp are serialized, so the cost scales linearly with the number of unique addresses the warp reads. The constant cache is at its best when threads in the same warp read only a few distinct locations, and when all of them read the same location, the guide says constant memory can be as fast as a register access. The compiler may place kernel parameters there too, which lets them be cached separately from the L1 data cache.
Underneath all of it sit caches nobody declares. Each SM has an L1 that is part of its unified data cache and a separate constant cache; a larger L2 is shared by every SM on the device. The best practices guide orders the whole hierarchy by access latency: global, local, and texture memory have the greatest access latency, followed by constant memory, then shared memory, then the register file. That ordering is what kernel design actually optimizes against. Every technique in the rest of this book, tiling most of all, is a way of paying the device memory trip once and then re-reading the value from somewhere further up this list.
A fourth tier: clusters#
One level sits above the thread block, and it is the level that qualifies the no-guarantee rule stated earlier. On GPUs of compute capability 9.0 and higher, a launch may group adjacent thread blocks into an optional tier called a cluster. Clusters are laid out in one, two, or three dimensions the way blocks and grids are, and specifying them changes neither the grid dimensions nor a block’s index within the grid. What changes is where the blocks run. All thread blocks in a cluster are executed in a single graphics processing cluster, the group of SMs the programming model uses to describe how a GPU is organized, and they are scheduled simultaneously. Because they are co-resident, threads in different blocks of the same cluster can communicate and synchronize with each other through Cooperative Groups. The maximum size of a cluster is hardware dependent and varies between devices.
Co-residency also makes a neighbor’s scratchpad reachable. Threads in a cluster can access the shared memory of every block in the cluster, which is called distributed shared memory: a thread block can read from, write to, and perform atomics in the shared memory of other thread blocks within its cluster. The Hopper tuning guide presents it as an intermediate step between the two options a kernel used to have, fitting the working set inside one block’s shared memory or falling back to global memory. An SM can use it simultaneously with L2 cache accesses, so code that moves data between SMs draws on the combined bandwidth of both. The access rules are the global memory rules rather than the shared memory ones: accesses to distributed shared memory should be coalesced and aligned to 32-byte segments where possible, and non-unit stride should be avoided.
Two limits bound the feature. The maximum portable cluster size supported is 8, and the H100 allows a nonportable cluster size of 16 for applications that opt in by setting the cudaFuncAttributeNonPortableClusterSizeAllowed function attribute. Larger clusters are not free, because using them may reduce the maximum number of active blocks across the GPU, which is why the guide tells applications launching cluster kernels to compute occupancy with cudaOccupancyMaxActiveClusters rather than with the block-level API. What Blackwell keeps of this and what it adds is in current hardware.