2.1Execution Model

This chapter exists to make the rest of the book readable: enough GPU programming to follow the architecture discussions — and essential if your research ever changes the hardware/software interface (researchers adding transactional memory to GPUs, for instance, had to write their own benchmarks because no GPU supports TM — see §5.3).

The central sleight of hand of GPU programming: the hardware is wide SIMD, but CUDA and OpenCL present a MIMD-like model. You launch a large array of scalar threads, each free to take its own control-flow path and touch arbitrary memory. At runtime, the hardware silently gangs these threads into warpswarpNVIDIA's group of 32 scalar threads executed in lockstep on SIMD hardware; the book's generic unit of SIMT execution.Glossary → (NVIDIA, 32 threads) or wavefrontswavefrontAMD's equivalent of a warp: 64 threads executed in lockstep.Glossary → (AMD, 64) that execute in lockstep on SIMD units, exploiting their regularity and spatial locality. That combination is the SIMTSIMTSingle-instruction, multiple-thread: the execution model GPU cores use to run thousands of scalar threads on SIMD hardware.Glossary → model. (One portability caveat the book flags, from Seo et al.: OpenCL code carefully tuned for a GPU can perform poorly on a CPU — the model travels, the performance doesn’t.)

From a loop to a grid of threads

Start with the book’s running example: SAXPY (single-precision a·x plus ySAXPYSingle-precision a·x + y (BLAS); the book's running example kernel.Glossary →, from the BLAS library — simple, useful for building things like Gaussian elimination, and a classic teaching example). Serial C first:

void saxpy_serial(int n, float a, float *x, float *y)
{
    for (int i = 0; i < n; ++i)
      y[i] = a*x[i] + y[i];
}
main() {
    float *x, *y;
    int n;
    // omitted: allocate CPU memory for x and y and initialize contents
    saxpy_serial(n, 2.0, x, y); // Invoke serial SAXPY kernel
    // omitted: use y on CPU, free memory pointed to by x and y
}

And the CUDA version — the for loop is gone, dissolved into threads:

__global__ void saxpy(int n, float a, float *x, float *y)
{
     int i = blockIdx.x*blockDim.x + threadIdx.x;
     if(i<n)
       y[i] = a*x[i] + y[i];
}
int main() {
    float *h_x, *h_y;
    int n;
    // omitted: allocate CPU memory for h_x and h_y and initialize contents
    float *d_x, *d_y;
    int nblocks = (n + 255) / 256;
    cudaMalloc( &d_x, n * sizeof(float) );
    cudaMalloc( &d_y, n * sizeof(float) );
    cudaMemcpy( d_x, h_x, n * sizeof(float), cudaMemcpyHostToDevice );
    cudaMemcpy( d_y, h_y, n * sizeof(float), cudaMemcpyHostToDevice );
    saxpy<<<nblocks, 256>>>(n, 2.0, d_x, d_y);
    cudaMemcpy( h_x, d_x, n * sizeof(float), cudaMemcpyDeviceToHost );
    // omitted: use h_y on CPU, free memory pointed to by h_x, h_y, d_x, and d_y
}

The __global__ keyword marks saxpy as a kernelkernelThe code a GPU-computing application launches to run across many GPU threads.Glossary → — the function every GPU thread runs. Each iteration of the original loop becomes one thread executing lines 3–5. Follow the whole host/device choreography step by step (the same launch flow you saw in §1.2, now tied to real code):

float *h_x, *h_y;
int n;
// allocate + initialize h_x, h_y

The thread hierarchy: grid → blocks → warps → threads

The launch syntax saxpy<<<nblocks, 256>>>(…) — CUDA’s kernel configuration syntaxkernel configuration syntaxCUDA's <<<blocks, threads>>> launch notation setting grid/CTA dimensions.Glossary → — creates a single gridgridThe full set of thread blocks launched by one kernel call.Glossary → of nblocks thread blocksthread block / CTACooperative thread array: a group of warps that shares a core's scratchpad and can barrier-synchronize; the unit assigned to a SIMT core.Glossary → (NVIDIA also says cooperative thread array, CTA), each containing 256 threads, which the hardware will slice into warps. Grids and blocks have x, y, and z dimensions; unspecified ones default to zero for every thread (as here). The kernel arguments (n, 2.0, d_x, d_y) are delivered to every thread.

Each thread finds its share of the data by computing a global index from three built-ins: blockIdx.x (which block am I in?), blockDim.x (how big is a block? — 256 here), and threadIdx.x (which thread am I within my block?). Watch blocks map onto the array one tick at a time (toy sizes: n = 20, blockDim = 8, warps of 4), and click any thread to compute its index — including the tail threads whose if (i<n) guard fails:

Note what the compiler and hardware are doing for you here: each thread appears to execute independently — the lockstep warps underneath are invisible at this level. (Chapter 3 is about what it costs to maintain that illusion.)

Memory: h_ and d_, and what the SAXPY example doesn’t show

Following NVIDIA convention, h_ prefixes pointers into CPU (host) memory and d_ into GPU (device) memory. cudaMalloc asks the driverdriverCPU-side software that conveys the kernel, thread count and data locations to the GPU and signals it to start.Glossary → to allocate device memory; cudaMemcpy asks it to copy — explicitly, in the general separate-memories case the book (and this site) focuses on. NVIDIA’s Unified Memoryunified memorySoftware/hardware support (NVIDIA Pascal+) that automatically migrates data between CPU and GPU memory via virtual memory.Glossary → can instead migrate data transparently in both directions, the runtime and hardware copying on your behalf; and integrated GPUs skip the copies entirely. Datacenter GPUs for machine learning still overwhelmingly use their own DRAM.

Beyond SAXPY, three programming-model facts matter for later chapters:

1. Fast intra-CTA communication. Threads of a CTA share a per-core scratchpadscratchpad memoryFast per-core, software-managed memory through which threads on a core communicate.Glossary →: NVIDIA’s shared memoryshared memoryNVIDIA's name for the per-SM scratchpad (16–64 KB), declared with __shared__; a software-controlled cache.Glossary → (one per SM, 16–64 KB, split among the CTAs resident there; declared with __shared__), AMD’s local data storelocal data store (LDS)AMD GCN's per-CU scratchpad memory, also used to pass parameters between graphics shaders.Glossary → (LDS) — plus, uniquely on AMD, a GPU-wide global data storeglobal data store (GDS)AMD GCN's GPU-wide scratchpad shared by all compute units.Glossary → (GDS). A scratchpad is a software-controlled cache: it wins when you can identify data that’s reused often and predictably (in graphics, e.g., passing parameters between vertex and pixel shaders). Hardware caches exist too, but miss unpredictably.
2. Fast intra-CTA synchronization. Threads in the same CTA synchronize with hardware barrierbarrierHardware-supported instruction letting all threads of a CTA synchronize cheaply.Glossary → instructions. Threads in different CTAs can only communicate through global memoryglobal memoryThe address space accessible to all threads; slower and more energy-costly than shared memory; the only way threads in different CTAs communicate.Glossary → — more expensive in both time and energy.
3. Nested parallelism. CUDA Dynamic ParallelismCUDA Dynamic Parallelism (CDP)Kepler-era feature letting kernels launch kernels, targeting load imbalance in irregular applications.Glossary → (Kepler) lets a kernel launch further work from the GPU itself, because irregular, data-dependent applications otherwise leave the hardware underutilized through load imbalance — the same disease Dynamic Warp Formation research attacks in §3.4.