GPU Kernels and Warps

A GPU kernel is a function run by thousands of threads; hardware executes those threads in lockstep groups called warps.

The execution model

A kernel is a function launched to run across a grid of threads. The grid is divided into thread blocks, and each block is scheduled onto one streaming multiprocessor (SM). The programmer writes code from the point of view of a single thread; the hardware replicates it across the whole grid, with each thread using its index to select the data it works on. This is the single-instruction, multiple-thread (SIMT) model.

Warps

Inside a block, threads are grouped into warps of 32 (on current architectures), and a warp is the unit the hardware actually schedules. All 32 threads in a warp share one program counter and execute the same instruction each cycle on their own data. The SM holds many warps resident at once and switches between them with zero overhead: when one warp stalls waiting on memory, the scheduler issues from another. This latency hiding, not fast individual threads, is the source of GPU throughput.

• Grid to blocks to warps to threads is the launch hierarchy.
• A warp is 32 threads executing one instruction in lockstep.
• Blocks map to one SM and can cooperate via shared memory.
• Many resident warps hide memory latency by fast context switching.

A minimal kernel

# CUDA-style pseudocode
# __global__ void axpy(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];
# }
# launch: axpy<<<(n+255)/256, 256>>>(n, 2.0f, x, y);

Why the model matters for performance

Because the warp is the scheduling unit, two effects dominate GPU efficiency: threads in a warp that take different branches (warp divergence) run serially, and threads in a warp that touch scattered memory addresses cannot be serviced in one transaction (poor coalescing). Writing kernels so a warp stays on one code path and touches contiguous memory is the foundation of GPU performance tuning, explored in the pages that follow.

Particle-tracking and field-update kernels in fusion codes are natural GPU targets when the per-thread work is uniform and memory access is regular.