Warp and SM: the GPU's real unit of execution

In the programming model you write threads, but the hardware does not schedule one thread at a time. 32 consecutive threads are bound into a warp, share one program counter, and execute the same instruction in lockstep. That is NVIDIA's SIMT (Single Instruction, Multiple Thread). Understanding warps is the line between writing CUDA that is correct and writing CUDA that is fast.

┌──────────────────────── SM ─────────────────────────┐
│  Warp Scheduler ×4   each cycle picks a ready warp  │
│         │                                           │
│  ┌──────┴──────┬───────────┬──────────┐             │
│  │ CUDA Core   │  LD/ST    │  SFU     │  Tensor     │
│  │  ×128       │  units    │          │  Core ×4    │
│  └─────────────┴───────────┴──────────┴────────────┘│
│                                                     │
│  Register File   256 KB (SM-wide, split per thread) │
│  Shared Memory / L1   up to 164 KB (configurable)   │
└─────────────────────────────────────────────────────┘
              │
        ┌─────┴─────┐
        │  L2 Cache │  chip-wide, tens of MB
        └─────┬─────┘
              │
        ┌─────┴─────┐
        │   HBM     │  device memory, tens of GB, 1-3 TB/s
        └───────────┘

How a block actually runs

1. 1.The scheduler assigns a whole block atomically to one SM. A block is never split across two SMs.
2. 2.The SM slices the block into warps by linear tid. blockDim = 256 is 8 warps.
3. 3.If blockDim is not a multiple of 32, the last warp is filled out but some threads stay inactive. Those slots waste issue bandwidth.
4. 4.Each cycle, the warp scheduler picks a ready warp among those resident on the SM and issues an instruction. When one warp is waiting on memory, the scheduler switches to another. The switch costs nothing, because every warp's registers stay physically allocated.

Warp divergence: what an if really costs

A warp has one program counter. When threads in a warp take different sides of an if, the hardware cannot run both paths at once. It serializes them: first the then path (threads that took else are masked off and idle), then the else path (the mask flips). The two paths add, and that sum is warp divergence.

1__global__ void bad(float* x) {2    int i = blockIdx.x * blockDim.x + threadIdx.x;3 4    if (i % 2 == 0) {         // half the warp each way: 100% divergence5        x[i] = expensiveA(x[i]);6    } else {7        x[i] = expensiveB(x[i]);8    }9    // Wall time ≈ expensiveA + expensiveB, not one or the other10}

1__global__ void good(float* x) {2    int i = blockIdx.x * blockDim.x + threadIdx.x;3 4    if ((i / warpSize) % 2 == 0) {   // every thread in the warp takes the same path5        x[i] = expensiveA(x[i]);6    } else {7        x[i] = expensiveB(x[i]);8    }9    // Each warp runs one path. Cost falls back to a true either-or10}

Synchronization: __syncthreads and __syncwarp

1__global__ void stencil(const float* in, float* out, int n) {2    __shared__ float tile[BLOCK + 2 * RADIUS];3 4    int gid = blockIdx.x * blockDim.x + threadIdx.x;5    int lid = threadIdx.x + RADIUS;6 7    tile[lid] = in[gid];                       // each thread writes its own cell8    if (threadIdx.x < RADIUS) {                // the first few threads also copy the halo9        tile[lid - RADIUS]     = in[gid - RADIUS];10        tile[lid + blockDim.x] = in[gid + blockDim.x];11    }12 13    __syncthreads();     // barrier: nobody reads until the whole tile is written14 15    float sum = 0.0f;16    for (int d = -RADIUS; d <= RADIUS; ++d) sum += tile[lid + d];17    out[gid] = sum;18}

Starting with Volta (sm_70), NVIDIA added Independent Thread Scheduling: each thread in a warp has its own program counter, and after divergence they are no longer guaranteed to reconverge on their own. Consequence: old code that assumed "threads in a warp are implicitly in sync, so no barrier is needed" (especially warp-level reductions) is wrong after Volta. Every intra-warp data exchange now needs an explicit masked sync.

1// Wrong: Kepler-era pattern. After Volta this is no longer guaranteed2volatile float* v = sdata;3if (tid < 32) { v[tid] += v[tid + 32]; v[tid] += v[tid + 16]; /* ... */ }4 5// Correct: explicit sync6if (tid < 32) {7    float val = sdata[tid] + sdata[tid + 32];8    for (int offset = 16; offset > 0; offset >>= 1) {9        val += __shfl_down_sync(0xffffffff, val, offset);10    }11    if (tid == 0) out[blockIdx.x] = val;12}