CUDA Programming Model Explained: Grids, Blocks, Warps and Memory

The CUDA programming model runs one function, a kernel, across thousands or millions of GPU threads. Threads are grouped into thread blocks, blocks form a grid, and each block runs on one streaming multiprocessor (SM), where hardware executes it 32 threads at a time as warps. Performance depends mostly on how those threads touch memory.

That paragraph is the whole model. The rest of this guide is the detail where the speed lives: why a block of 100 threads wastes more than a fifth of its lanes, why reading a column of a 32 by 32 shared memory tile is a 32-way bank conflict, and why a reduction that worked for a decade can break on newer GPUs. We go from host and device through warps and every memory space to two hands-on kernels.

Key takeaways

• A kernel launch creates a grid of thread blocks. All threads of a block run on one SM and can share on-chip shared memory and synchronise with __syncthreads(). Blocks in a grid cannot rely on each other, because CUDA makes no promise about the order or concurrency in which they run.
• An SM splits each block into warps of 32 threads. When threads in a warp take different branches, the warp runs each path in turn with the other lanes masked off.
• From compute capability 7.0 (Volta), independent thread scheduling gives each thread its own program counter and call stack. Code that assumed a warp moves in lockstep must now use __syncwarp() or the _sync intrinsics.
• Thread block clusters (compute capability 9.0 and up) group blocks that are co-scheduled on one GPC and can use each other's shared memory.
• The memory hierarchy runs from per-thread registers, through per-block shared memory carved from the same storage as L1, to L2 and global memory. "Local" memory is per-thread in scope but lives in device memory.
• Global memory is served in 32-byte sectors. A warp reading 32 consecutive 4-byte words needs 4 sectors and uses 100 percent of what it fetches. A warp reading words 32 bytes or more apart needs 32 sectors and uses 12.5 percent.
• Shared memory has 32 banks of 4-byte words. A stride of s words between lanes gives a conflict degree equal to the greatest common divisor of s and 32, which is why padding a 32 by 32 tile to 32 by 33 removes a 32-way conflict.
• Compute capability decides which features exist: hardware asynchronous copies from 8.0, clusters, distributed shared memory and the Tensor Memory Accelerator from 9.0, and architecture-specific (a) and family-specific (f) compiler targets from 9.0 and 10.0.

A note on method before we start. Every specification below comes from NVIDIA's own documentation: the CUDA Programming Guide (version 13.4.2, last updated 10 September 2026), the CUDA C++ Best Practices Guide, the Hopper Tuning Guide, the NVCC documentation, the CUDA Toolkit 13.4 Update 2 release notes, and the Volta and Ampere architecture whitepapers. Every computed number was checked in Python. The code samples follow NVIDIA's documented APIs but were not compiled or run, because this article was written on a Mac with no NVIDIA GPU. The only measured performance figures are NVIDIA's own, labelled with the GPU they came from.

If you want the hardware view first, our explainer on AI chip architecture and the guide to what an AI accelerator is and how it works cover the silicon. This post is about the programming model that sits on top of it.

Host and device: the machine CUDA assumes

CUDA assumes a heterogeneous system. The Programming Guide calls the CPU and its directly attached memory the host and host memory, and a GPU with its directly attached memory the device and device memory. On a system on chip the two can share one physical memory, but the model still distinguishes them.

A CUDA application always starts on the CPU. Host code uses CUDA APIs to allocate device memory, copy data across, launch work on the GPU and wait for it. A function that runs on the GPU is called a kernel, "for historical reasons" in the guide's words, and starting one is launching it. The CPU and GPU can run at the same time, and the guide notes that the best performance usually comes from keeping both busy.

            HOST                                      DEVICE
 +---------------------------+            +-------------------------------------+
 |  CPU cores                |            |  GPC   GPC   GPC   GPC   ...        |
 |  host memory (system RAM) | <--------> |  [SM][SM][SM] [SM][SM][SM] ...      |
 |                           |  PCIe or   |                                     |
 |  main(), cudaMalloc,      |  NVLink    |  L2 cache (shared by all SMs)       |
 |  cudaMemcpy, launch       |            |  device memory (global, DRAM/HBM)   |
 +---------------------------+            +-------------------------------------+

On all currently supported systems the CPU and GPU share one unified virtual address space, but that does not make every pointer usable everywhere. Host code cannot dereference memory from cudaMalloc, and on many systems a kernel cannot touch memory from malloc. Unified memory, covered later, relaxes this.

The guide's hardware model is deliberately simple. A GPU is a set of SMs grouped into Graphics Processing Clusters (GPCs). Each SM has a register file, a unified data cache and functional units. The guide adds a caution worth repeating: the real hardware layout may differ, and those differences "do not affect correctness" of code written to the model. The model is a contract about behaviour, not a floor plan.

Kernels and the launch configuration

You mark a kernel with __global__, give it a void return type and launch it with an execution configuration in triple angle brackets. The full form takes four values.

// kernel<<<gridDim, blockDim, dynamicSharedBytes, stream>>>(args...);
dim3 block(256);                         // 256 threads per block, 1D
dim3 grid((n + block.x - 1) / block.x);  // enough blocks to cover n elements
myKernel<<<grid, block, 0, 0>>>(d_in, d_out, n);

The first two values are the grid and block dimensions. The third is the number of bytes of dynamically sized shared memory each block gets. The fourth is the stream the launch is queued in. The last two are optional and default to zero and the default stream. The guide also documents cudaLaunchKernelEx, which takes a configuration struct and is one way to pass launch attributes such as a cluster size.

Inside the kernel, four built-in variables give each thread its identity. The guide defines them as follows.

The standard one-dimensional global index is threadIdx.x + blockIdx.x * blockDim.x. For two dimensions you compute a row and a column the same way from the y and x components. The guide stresses that multi-dimensional blocks and grids are "for convenience only and does not affect performance". Threads are linearised with x fastest, then y, then z, so threadIdx.y has a stride of blockDim.x in the linear thread ID. That linear ID is what decides which warp a thread lands in.

The limits you will hit

The Programming Guide's Table 30 gives the launch limits for every currently supported compute capability, 7.5 through 12.x. Most are the same everywhere.

A one-dimensional grid can therefore address up to 2,147,483,647 blocks of 1,024 threads, which is about 2.2 trillion threads in a single launch. You will run out of memory long before you run out of index space. The limit people actually hit is the 1,024 threads per block, which applies to the product of the three block dimensions: a 32 by 32 block is legal, a 32 by 64 block is not.

How a grid maps onto the GPU

When a kernel launches, a hardware scheduler hands thread blocks to SMs. The guide is explicit about what you can and cannot rely on.

• All threads of one block run on a single SM and, in most cases, run to completion there. The exception the guide notes is CUDA Dynamic Parallelism, where a block can be suspended to memory.
• There is no guarantee about which SM gets which block, or in what order blocks run. Programs "cannot rely on a specific scheduling order or scheme for correct execution".
• An SM keeps accepting blocks while it has free resources: threads, warps, registers, shared memory and block slots. When no SM can take another block, the scheduler waits for running blocks to finish.

This is what lets one compiled kernel run on a GPU with a handful of SMs or more than a hundred. The price is that blocks must not depend on each other within a launch. If block 7 spins on a flag that block 9 sets, and block 9 never gets an SM, the kernel hangs. Cooperative launches, covered later, are the sanctioned way around this.

Waves: a worked example on A100

The NVIDIA Ampere architecture whitepaper gives the A100 Tensor Core GPU 108 SMs. The Programming Guide gives compute capability 8.0 a maximum of 2,048 resident threads and 32 resident blocks per SM. Take a vector of 2^24 elements (16,777,216) with 256 threads per block.

blocks needed          = ceil(16,777,216 / 256)          = 65,536
blocks resident per SM = 2,048 threads / 256 per block   = 8
resident blocks (GPU)  = 108 SMs x 8                      = 864
resident threads (GPU) = 108 x 2,048                      = 221,184
"waves" of blocks      = 65,536 / 864                     = 75.85

That assumes registers and shared memory do not limit residency first. In the 76th wave only about 85 percent of the SMs have work, and a kernel with only two or three waves loses a visible share of its time to such a tail. This is one reason the grid-stride loop, where you launch a fixed number of blocks sized to the GPU and each thread loops over several elements, is a common pattern. We use it in the reduction later.

Warps and SIMT

Within a block, the guide says, threads are organised into groups of 32 called warps, and each warp executes in a Single-Instruction, Multiple-Thread (SIMT) fashion. Each thread in a warp gets a lane number from 0 to 31. Warps are formed from consecutive linear thread IDs, so threads 0 to 31 are warp 0, threads 32 to 63 are warp 1, and so on. The number of warps in a block is the thread count divided by 32, rounded up.

Block of 128 threads (blockDim.x = 128)

 thread IDs:  0 ......... 31 | 32 ........ 63 | 64 ........ 95 | 96 ....... 127
 warps:       [   warp 0    ] [    warp 1    ] [    warp 2    ] [   warp 3    ]
 lanes:        0 ......... 31   0 ......... 31   0 ......... 31   0 ......... 31

Block of 100 threads: ceil(100/32) = 4 warps, 128 lanes, 28 lanes idle
                       idle share of the block's lanes = 28/128 = 21.9%

That is why the guide recommends block sizes that are multiples of 32. Any size is legal, but the last warp of an odd-sized block carries dead lanes for its whole life, which wastes functional units and makes memory accesses less efficient.

How SIMT differs from SIMD

The guide draws the line clearly. SIMD vector hardware exposes its width to software: you write code for an 8-wide or 16-wide register and handle the edge cases yourself. SIMT instructions describe the behaviour of a single thread. You write scalar code for one thread, and the hardware groups 32 of them. Each thread can follow its own control flow, so SIMT has no fixed data width from the programmer's point of view.

For correctness, the guide says, you can essentially ignore SIMT. For performance you cannot, because the warp is the unit that issues instructions and the unit that issues memory requests. The two biggest consequences are branch divergence and memory coalescing.

Branch divergence

A warp executes one common instruction at a time. When threads of the same warp take different paths through a data-dependent branch, the warp executes each path that any of its threads took, disabling the lanes not on that path. Different warps never diverge from each other in this sense, since they are scheduled independently.

if (threadIdx.x % 2 == 0) {
    a[i] = f(a[i]);   // path A: even lanes
} else {
    a[i] = g(a[i]);   // path B: odd lanes
}

Warp timeline for the branch above (32 lanes, X = active, . = masked)

 path A:  X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.
 path B:  .X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X.X
 => both paths issued; lane utilisation across the two = 50%

If you change the condition to (threadIdx.x / 32) % 2 == 0, whole warps go one way or the other and no warp diverges. The amount of work is the same, but each warp issues only one path. This exact transformation is the first optimisation in the reduction walkthrough below.

Divergence is not always a disaster. A branch that is uniform across a warp costs nothing extra. The cost is in long divergent regions and in loops whose trip count varies by lane, where the warp waits for its slowest thread.

Independent thread scheduling from Volta onward

Before compute capability 7.0, a warp had a single program counter shared by all 32 threads, plus an active mask. The Programming Guide notes a consequence: threads of the same warp in divergent regions could not signal each other or exchange data, and algorithms that use fine-grained locks could deadlock depending on which warp the contending threads came from.

The Volta architecture whitepaper (NVIDIA, August 2017) introduced independent thread scheduling. The GPU keeps execution state per thread, including a program counter and call stack, and can yield at per-thread granularity. In the Programming Guide's description, a schedule optimiser groups active threads from the same warp into SIMT units, so throughput stays high, but threads "can now diverge and reconverge at sub-warp granularity".

Pre-Volta (one PC per warp)            Volta and later (PC and stack per thread)

 warp PC ---> A ---> B ---> join        lane 0 PC ---> A ----> X ---> ...
                                        lane 1 PC ---> B --> Y ---> ...
 all lanes reconverge at a fixed point  scheduler interleaves paths; lanes may
                                        not reconverge where you expect

The catch is warp-synchronous code. For years, CUDA programmers wrote code that assumed the 32 threads of a warp move in lockstep at every instruction, and so skipped synchronisation inside a warp. The guide says such code "should be revisited", naming synchronisation-free intra-warp reductions as the example, and says to use __syncwarp() to make the intent explicit. The intrinsics that exchange data within a warp, such as __shfl_sync, __shfl_down_sync, __ballot_sync and, on compute capability 8.0 and later, __reduce_add_sync, all take an explicit mask of participating lanes for the same reason. The guide itself suggests the CUB warp-wide primitives for warp reductions where possible. We will see a concrete broken example in the reduction section.

Thread block clusters on Hopper and later

Compute capability 9.0, the Hopper generation, added an optional level between block and grid. A thread block cluster is a group of blocks laid out in one, two or three dimensions. The Programming Guide describes three properties.

1. Specifying clusters does not change the grid dimensions or block indices. A block simply also has a position within its cluster.
2. All blocks in a cluster are scheduled at the same time on SMs within a single GPC.
3. Threads in a cluster can read, write and perform atomics on the shared memory of every block in the cluster. That pooled space is distributed shared memory.

Grid of 8 x 2 blocks with 2 x 1 clusters

 cluster(0,0)  cluster(1,0)  cluster(2,0)  cluster(3,0)
 [B0,0][B1,0]  [B2,0][B3,0]  [B4,0][B5,0]  [B6,0][B7,0]
 [B0,1][B1,1]  [B2,1][B3,1]  [B4,1][B5,1]  [B6,1][B7,1]
 cluster(0,1)  cluster(1,1)  cluster(2,1)  cluster(3,1)

 Each cluster: both blocks co-scheduled on SMs of one GPC;
 block 0 can load from block 1's __shared__ array and vice versa.

The Hopper Tuning Guide gives the sizing rules. The maximum portable cluster size is 8 blocks. The H100 GPU allows a non-portable size of 16 if the kernel opts in with the cudaFuncAttributeNonPortableClusterSizeAllowed attribute, and larger clusters can reduce how many blocks are active at once. The tuning guide recommends computing occupancy for cluster kernels with cudaOccupancyMaxActiveClusters. It also notes that distributed shared memory traffic can run alongside L2 traffic, so an SM can draw on both at once, and that accesses to distributed shared memory should be coalesced just like global ones.

Here is a minimal cluster kernel using the documented cooperative groups API. Blocks in a pair swap the contents of their shared arrays.

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

// Compile-time cluster shape of 2 x 1 x 1 blocks (requires CC 9.0 or later)
__global__ void __cluster_dims__(2, 1, 1) swap_halves(float* out)
{
    __shared__ float tile[256];               // assumes blockDim.x == 256
    cg::cluster_group cluster = cg::this_cluster();
    unsigned int rank = cluster.block_rank(); // 0 or 1 within the cluster

    tile[threadIdx.x] = static_cast<float>(rank * 1000 + threadIdx.x);
    cluster.sync();                           // all blocks started and wrote

    // Pointer into the *other* block's shared memory (distributed shared memory)
    float* peer = cluster.map_shared_rank(tile, rank ^ 1);
    out[blockIdx.x * blockDim.x + threadIdx.x] = peer[threadIdx.x];

    cluster.sync();                           // peer must not exit while we read
}

The second cluster.sync() matters. The Programming Guide warns that all distributed shared memory operations must finish before a block exits, because once a block is gone its shared memory is gone with it.

The memory hierarchy

CUDA exposes several memory spaces. They differ in where they physically live, who can see them and how long they last. The table below combines the Programming Guide's descriptions of each space.

             per thread                per block / cluster            per device
        +----------------+         +---------------------------+   +-------------+
 thread | registers      |         | shared memory   | L1      |   |  L2 cache   |
        | (up to 255 x   |  <--->  | (carve-out of the unified |<->|  (all SMs)  |<-> device DRAM
        |  32-bit)       |         |  data cache, 32 banks)    |   +-------------+    (global, local,
        +----------------+         +---------------------------+                       constant, texture)
          fastest, smallest            on chip, per SM                 on chip          off chip, largest

Registers

Registers live on the SM and belong to one thread. The compiler allocates them. Table 31 of the Programming Guide gives every supported compute capability 64K 32-bit registers per SM, at most 64K per block and at most 255 per thread. You can cap registers per thread with the -maxrregcount option to NVCC or per kernel with __launch_bounds__. Capping can let more blocks fit on an SM, but the guide warns that it can also cause register spilling: values the compiler cannot keep in registers are written out to local memory and read back later.

To see what the compiler did, NVCC's --resource-usage option prints registers and shared memory per kernel, which the guide recommends as the first step of any occupancy calculation.

Local memory

Local memory is the most misnamed space in CUDA. The guide states that it is thread-local in scope but "the physical location of local memory is in the global memory space". The compiler is likely to place three kinds of variable there:

1. arrays it cannot prove are indexed with constant quantities,
2. large structures or arrays that would use too many registers,
3. anything at all, once the kernel runs out of registers and spills.

Because it lives in device memory, local memory has global memory's latency and bandwidth. It is laid out so that consecutive 32-bit words belong to consecutive thread IDs, which means accesses are fully coalesced when every thread in a warp touches the same relative address, such as the same index of the same local array. The limit is 512 KB of local memory per thread on all current compute capabilities. A small private array indexed by a loop variable that the compiler fully unrolls usually stays in registers. The same array indexed by a data-dependent value usually goes to local memory.

Shared memory and L1

Shared memory is accessible by all threads of a block and lives on the SM in the same physical structure as the L1 cache, which NVIDIA calls the unified data cache. The guide calls it "a user-managed scratchpad". It has much higher bandwidth and lower latency than global memory, but it is small and it is allocated per block: every resident block on an SM takes its own slice.

You declare it statically, with a fixed size, or dynamically, with extern __shared__ and the third launch parameter.

__global__ void staticSmem(const float* in) {
    __shared__ float tile[32][33];  // size fixed at compile time
    // ...
}

__global__ void dynamicSmem(const float* in) {
    extern __shared__ float buf[];  // size given at launch
    // ...
}
// dynamicSmem<<<grid, block, block.x * sizeof(float)>>>(d_in);

The split between L1 and shared memory is configurable. The Programming Guide lists the unified data cache size and the shared memory capacities you can select per compute capability. A kernel can state a preference with cudaFuncSetAttribute and cudaFuncAttributePreferredSharedMemoryCarveout, and must opt in explicitly with cudaFuncAttributeMaxDynamicSharedMemorySize to use more than 48 KB of dynamic shared memory per block.

Source: CUDA Programming Guide 13.4.2, Tables 31 and 32. KB here means 1,024 bytes. The Hopper Tuning Guide puts the jump from compute capability 8.0 to 9.0 at 39 percent (164 KB to 228 KB per SM), which checks out: 228 divided by 164 is 1.39. On every row, the per-block maximum is 1 KB below the per-SM maximum.

L2 cache

The L2 cache sits on the GPU and is shared by every SM. Its size varies widely by product, so the guide tells you to query l2CacheSize from cudaGetDeviceProperties rather than assume. For scale, the Volta whitepaper lists 6,144 KB of L2 for a full GV100, the Ampere whitepaper lists 40 MB for A100, and the Hopper Tuning Guide lists 50 MB for H100. From compute capability 8.0 you can mark a window of global memory as persisting in L2, which the guide covers under L2 cache control. That helps when a small lookup table is hit repeatedly while a large stream of data flows past it.

Global memory

Global memory, also called device memory, is the main store. The guide compares it to RAM on a CPU. Allocations persist until freed with cudaFree or until the application exits. Since kernels launched from the host return void, writing to global memory is the only way a kernel can hand numbers back. Every thread in the grid can see it, so you are responsible for avoiding races, either by partitioning the data so each element has one writer or by using atomics.

Constant memory

Constant memory is declared with __constant__ at file scope, is read-only inside kernels, and is set from the host with cudaMemcpyToSymbol. The guide gives 64 KB of constant memory per device and an 8 KB constant cache working set per SM on every current compute capability. It works best when all threads in a warp read the same address, such as a filter coefficient, since the value is then served once and broadcast. The Best Practices Guide notes that when threads in a warp read different constant addresses, the reads are serialised, so constant memory is the wrong home for a per-thread lookup table.

__constant__ float coeffs[4];

__global__ void poly(const float* x, float* y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float v = x[i];
        y[i] = ((coeffs[3] * v + coeffs[2]) * v + coeffs[1]) * v + coeffs[0];
    }
}
// host: cudaMemcpyToSymbol(coeffs, h_coeffs, sizeof(h_coeffs));

Texture and surface memory

Older CUDA code often reads arrays through textures because that helped on early GPUs. The current guide is blunt: texture and surface memory "do not provide any performance advantage for non-graphics applications in CUDA on any currently supported NVIDIA GPU". The APIs remain for real texture work, such as rendering with OptiX.

Coalescing: worked address patterns

The Programming Guide states the rule in one sentence: global memory is accessed through 32-byte memory transactions. When a warp executes a load, the hardware gathers the 32 addresses, works out which 32-byte sectors they fall in, and fetches each needed sector once. Efficiency is the ratio of bytes the threads asked for to bytes actually moved.

The best way to build intuition is to compute sector counts for real address patterns. All addresses below are byte offsets from a 32-byte aligned base, for one warp of 32 threads, lane i from 0 to 31.

The first and fourth rows are the two examples the guide itself works through: 4 transactions for perfect coalescing, and 32 transactions moving 1,024 bytes for 128 useful bytes in the worst case, "only 12.5%". The misaligned row shows why allocation alignment matters: one element of offset adds a fifth sector and costs 20 percent. The Best Practices Guide states that memory from cudaMalloc is aligned to at least 256 bytes, so this usually appears when you index a sub-array at an odd offset.

The guide also stresses something people forget: there is no requirement that lane i read element i. Any permutation of addresses within the same sectors coalesces just as well. What matters is the set of sectors, not the order.

Where strided access comes from

Strided access usually appears in two places.

Array of structures. If each element is a struct of eight floats and a kernel only reads field x, consecutive lanes read addresses 32 bytes apart: the 12.5 percent row. Rewriting the data as a structure of arrays, one array per field, turns that into the 100 percent row.

// Array of structures: lane i reads particles[i].x, stride 32 bytes
struct Particle { float x, y, z, vx, vy, vz, mass, charge; };
__global__ void readX_aos(const Particle* p, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = p[i].x;      // 32 sectors per warp
}

// Structure of arrays: lane i reads x[i], stride 4 bytes
__global__ void readX_soa(const float* x, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = x[i];        // 4 sectors per warp
}

Column access in row-major matrices. A naive matrix transpose reads rows and writes columns, or the other way round. Whichever side walks a column of a row-major matrix of width m has lanes m * 4 bytes apart, which for any useful matrix is far more than 32 bytes. The guide's transpose example fixes this by staging a 32 by 32 tile in shared memory: read a tile with coalesced row reads, then write it out with coalesced row writes of the transposed tile. That moves the strided access into shared memory, where the rules are different, and introduces the next topic.

Shared memory bank conflicts: worked

The guide describes shared memory as 32 banks, with successive 32-bit words mapped to successive banks, each bank delivering 32 bits per clock. For a 4-byte word at byte address a within shared memory:

bank(a) = (a / 4) mod 32

If two or more lanes of a warp access different words in the same bank, the accesses to that bank are serialised. The guide names two exceptions: if several lanes read the same word, it is broadcast, and if several lanes write the same word, only one write happens and which lane wins is undefined.

For a warp where lane i reads word s * i, the number of lanes that land on the busiest bank is the greatest common divisor of s and 32. We checked that by brute force over all 32 lanes.

The stride 1, 2 and 3 rows match the guide's Figure 15. The rule of thumb that falls out is simple: an odd stride never conflicts with 32 banks, and a power-of-two stride conflicts by its own size, up to 32.

The transpose tile, padded

The guide's transpose example uses __shared__ float smemArray[32][32] with a 32 by 32 thread block. A warp has a fixed threadIdx.y and lanes threadIdx.x from 0 to 31.

• Writing smemArray[threadIdx.x][threadIdx.y] walks down a column. Consecutive lanes are 32 words apart, the stride 32 row of the table: a 32-way conflict.
• Reading smemArray[threadIdx.y][threadIdx.x] walks along a row: stride 1, no conflict.

The fix is to pad the inner dimension by one.

#define TILE 32

__global__ void transposeTiled(const float* a, float* c, int m)
{
    __shared__ float tile[TILE][TILE + 1];    // +1 column of padding

    int col = blockIdx.x * TILE + threadIdx.x;
    int row = blockIdx.y * TILE + threadIdx.y;
    if (row < m && col < m)
        tile[threadIdx.y][threadIdx.x] = a[row * m + col];  // coalesced read

    __syncthreads();

    int tcol = blockIdx.y * TILE + threadIdx.x;  // swap block coordinates
    int trow = blockIdx.x * TILE + threadIdx.y;
    if (trow < m && tcol < m)
        c[trow * m + tcol] = tile[threadIdx.x][threadIdx.y]; // coalesced write,
                                                             // column read of tile
}

With 33 words per row, walking down a column moves 33 words per step, and 33 is odd, so all 32 lanes land on different banks. The padding costs 128 bytes per tile, from 4,096 bytes to 4,224 bytes, which is a trivial share of a modern SM's shared memory.

Bank of element [row][0] for the first 4 rows

            32 x 32 tile              32 x 33 tile
 row 0:     bank  0                   bank  0
 row 1:     bank  0  (32 mod 32)      bank  1  (33 mod 32)
 row 2:     bank  0  (64 mod 32)      bank  2  (66 mod 32)
 row 3:     bank  0  (96 mod 32)      bank  3  (99 mod 32)
 => column walk: 32-way conflict      => column walk: conflict free

Padding is not the only fix. Swizzling, where the column index is XORed with a function of the row, avoids conflicts without wasting space and is what many high-performance libraries do. Padding is the version to learn first because it is easy to check by hand.

Synchronisation

CUDA gives you synchronisation at several scopes, and choosing the narrowest one that is correct is half of writing fast kernels.

Within a block: __syncthreads()

__syncthreads() blocks every thread in the block until all of them reach it, and guarantees that shared and global memory writes made before it are visible to all threads in the block after it. The canonical pattern is write, barrier, read.

__global__ void blockSum(const int* in, int* out)
{
    __shared__ int s[128];                     // assumes blockDim.x == 128
    s[threadIdx.x] = in[blockIdx.x * blockDim.x + threadIdx.x];
    __syncthreads();                           // all writes to s are visible
    if (threadIdx.x == 0) {
        int sum = 0;
        for (int i = 0; i < blockDim.x; ++i) sum += s[i];
        out[blockIdx.x] = sum;
    }
}

The rule that bites people: every thread of the block must reach the same __syncthreads(). Putting one inside a branch that only some threads take is undefined behaviour and in practice often hangs the block. If you need a barrier inside a conditional, restructure so the condition only guards the work and the barrier sits outside it.

Within a warp: __syncwarp() and the _sync intrinsics

__syncwarp(mask) synchronises the lanes named in the mask and orders their memory accesses. Called with no argument, it uses the full mask 0xFFFFFFFF. The warp-level data exchange intrinsics carry the same mask and let lanes read each other's registers directly, without going through shared memory. A warp-wide sum with shuffles looks like this.

__inline__ __device__ float warpSum(float v)
{
    // 5 steps for 32 lanes: offsets 16, 8, 4, 2, 1
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xFFFFFFFF, v, offset);
    return v;   // lane 0 holds the warp total
}

On compute capability 8.0 and later, the Programming Guide's feature table adds warp reduce functions such as __reduce_add_sync, which do the same for signed and unsigned integers in one call.

Across the device: atomics

There is no barrier across a whole grid in an ordinary launch. What you do have are atomic operations, which perform a read-modify-write on one memory location without interference from other threads. CUDA offers the classic functions such as atomicAdd and atomicCAS, and, through the CUDA C++ standard library, cuda::std::atomic and cuda::atomic_ref, which follow C++ semantics and let you choose a thread scope: block, device or system. A narrower scope can be cheaper because the hardware does not need to make the update visible as widely.

#include <cuda/atomic>

__global__ void sumAll(const float* x, float* total, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        cuda::atomic_ref<float, cuda::thread_scope_device> t(*total);
        t.fetch_add(x[i], cuda::std::memory_order_relaxed);
    }
}

That kernel is correct and slow. Sixteen million threads all hitting one address serialise on it. The fast pattern, which the reduction section builds, reduces within warps and blocks first and issues one atomic per block.

The feature table adds two atomic capabilities from compute capability 9.0: atomics on 128-bit integers in shared and global memory, and atomicAdd on float2 and float4 vectors in global memory.

Cooperative groups

Cooperative groups is a C++ library that makes groups of threads explicit objects you can pass around, partition and synchronise. The guide lists what it adds beyond the built-in barriers: groups smaller than a block, and groups larger than a block, up to the whole grid and even multiple GPUs.

#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;

__global__ void tiledSum(const int* in, int* out)
{
    cg::thread_block block = cg::this_thread_block();
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);

    int v = in[block.group_index().x * block.size() + block.thread_rank()];
    int warpTotal = cg::reduce(warp, v, cg::plus<int>());   // within one warp

    if (warp.thread_rank() == 0)
        atomicAdd(out, warpTotal);
}

For a grid-wide barrier, the guide says you must launch with cudaLaunchCooperativeKernel. That call is atomic in the guide's sense: if it succeeds, every requested block is resident on the device at once, which is what makes cg::this_grid().sync() safe. The trade-off is that the grid can be no larger than what fits on the GPU at once, so you size it with the occupancy API.

Asynchronous barriers

From compute capability 7.0, CUDA offers asynchronous barriers through cuda::barrier, which split a barrier into an arrive and a wait. A thread can signal that its part is done, do unrelated work, and only then wait for the others. The guide notes that compute capability 8.0 and later accelerate these barriers in hardware in shared memory and allow hardware-accelerated synchronisation of any subset of threads in a block. Earlier GPUs only accelerated whole-warp and whole-block barriers. Split barriers are the glue for the asynchronous copies in the next section.

Streams, events and asynchronous copies

Streams

A stream is an ordered queue of GPU work. Operations in one stream run in order. Operations in different streams may run concurrently or in any order relative to each other. CUDA has a default stream that receives anything launched without one. The guide explains that streams created with cudaStreamCreate are blocking with respect to the legacy default stream: work in the legacy default stream waits for them, and they wait for it. Streams created with cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking) do not synchronise with it. NVCC also offers a per-thread default stream, where each host thread gets its own default stream.

The classic use is overlapping copies with compute. Asynchronous host to device copies need page-locked (pinned) host memory, which you get from cudaMallocHost or cudaHostAlloc, or by registering an existing buffer with cudaHostRegister.

const int CHUNKS = 4;
size_t chunk = n / CHUNKS;              // assume n divisible by CHUNKS
float *h, *d;
cudaMallocHost(&h, n * sizeof(float));  // pinned, so copies can be async
cudaMalloc(&d, n * sizeof(float));

cudaStream_t s[CHUNKS];
for (int k = 0; k < CHUNKS; ++k) cudaStreamCreateWithFlags(&s[k], cudaStreamNonBlocking);

for (int k = 0; k < CHUNKS; ++k) {
    size_t off = k * chunk;
    cudaMemcpyAsync(d + off, h + off, chunk * sizeof(float), cudaMemcpyHostToDevice, s[k]);
    scale<<<(chunk + 255) / 256, 256, 0, s[k]>>>(d + off, chunk);
    cudaMemcpyAsync(h + off, d + off, chunk * sizeof(float), cudaMemcpyDeviceToHost, s[k]);
}
cudaDeviceSynchronize();

With four streams, the copy for chunk 1 can overlap the kernel for chunk 0. How much overlap you get depends on the copy engines and the interconnect, so measure it with a profiler.

Events

Events are markers you record into a stream. You can wait on them from the host, make another stream wait on them with cudaStreamWaitEvent, or measure elapsed time between two of them. Event timing measures GPU time in the stream, which avoids counting host overhead.

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start, stream);
myKernel<<<grid, block, 0, stream>>>(args);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);

float ms = 0.0f;
cudaEventElapsedTime(&ms, start, stop);

Asynchronous copies into shared memory

The traditional way to stage data in shared memory is a load into a register followed by a store to shared memory. That ties up registers and the thread. The Ampere whitepaper introduced an asynchronous copy instruction that loads data directly from global memory into shared memory, bypassing registers. The Programming Guide calls the instruction LDGSTS and documents it for compute capability 8.0 and later. It copies 4, 8 or 16 bytes per thread. Copies of 4 or 8 bytes go through L1. Copies of 16 bytes can bypass L1, so they do not pollute it. You use it through cuda::memcpy_async, cooperative_groups::memcpy_async or the lower-level __pipeline_memcpy_async, and you wait on it with a barrier or a pipeline.

#include <cooperative_groups.h>
#include <cooperative_groups/memcpy_async.h>
namespace cg = cooperative_groups;

__global__ void stageTile(const float* in, float* out)
{
    extern __shared__ float tile[];
    cg::thread_block block = cg::this_thread_block();
    size_t base = blockIdx.x * block.size();

    cg::memcpy_async(block, tile, in + base, sizeof(float) * block.size());
    // ... independent work can go here while the copy is in flight ...
    cg::wait(block);                       // copy complete and visible

    out[base + block.thread_rank()] = tile[block.thread_rank()] * 2.0f;
}

Compute capability 9.0 adds the Tensor Memory Accelerator (TMA). The Hopper Tuning Guide describes it as a more capable asynchronous copy engine: it moves one- to five-dimensional tensors between global and shared memory in both directions, and between the shared memories of SMs in the same cluster. A single thread can issue a large transfer and the rest of the block keeps working. That is the basis of the producer and consumer warp pipelines in modern matrix multiply kernels. TMA is a large topic in its own right, and the follow-up post in this series on the inside of the SM goes further into it.

Unified memory

Unified memory lets one allocation, called managed memory, be used from both CPU and GPU code. The runtime or the hardware moves or maps the data as needed. You create managed memory with cudaMallocManaged, with a memory pool of managed type, or with the __managed__ specifier. On systems with Heterogeneous Memory Management (HMM) or Address Translation Services (ATS), all system memory is managed implicitly.

The guide is careful to say that behaviour depends on operating system, kernel version, GPU and interconnect, and gives four paradigms, which you detect with three device attributes.

The attribute names are abbreviated in the header. The full names are cudaDevAttrConcurrentManagedAccess, cudaDevAttrPageableMemoryAccess and cudaDevAttrPageableMemoryAccessUsesHostPageTables.

With full support, the guide says managed memory is usually placed where it is first touched, migrated when another processor uses it, moved at page granularity under software coherence or cache-line granularity under hardware coherence, and may be oversubscribed: you can allocate more managed memory than the GPU physically has. On hardware with ATS, the CPU can access GPU-resident managed memory without migrating it, and the link supports native atomics.

float* x;
cudaMallocManaged(&x, n * sizeof(float));
for (size_t i = 0; i < n; ++i) x[i] = 1.0f;     // first touch on the CPU

int dev = 0;
cudaMemLocation loc{cudaMemLocationTypeDevice, dev};
cudaMemPrefetchAsync(x, n * sizeof(float), loc, 0, stream);  // move before use
scale<<<grid, block, 0, stream>>>(x, n);
cudaStreamSynchronize(stream);

The guide's performance advice is short: unified memory makes code correct with less effort, but "optimal performance is attained by keeping the migration of memory to a minimum". Without prefetching, the first kernel to touch a page pays for a fault and a migration. Prefetch hints and cudaMemAdvise let you tell the runtime where data should live. The prefetch call follows the signature in the current guide, with a cudaMemLocation and a flags argument. Earlier toolkits used a different signature, so check the API reference for the version you build with.

Compute capability: how features are gated

Every NVIDIA GPU has a compute capability, a major and minor version such as 8.0 or 9.0. It is not a product name or a speed rating. It identifies the feature set and the limits. You can read it with one command, which the guide documents.

nvidia-smi --query-gpu=name,compute_cap --format=csv

At run time, cudaDeviceGetAttribute with cudaDevAttrComputeCapabilityMajor and cudaDevAttrComputeCapabilityMinor gives the same answer.

The guide's Table 29 shows which features arrive where. These are the ones this article touches.

Resource limits also change with compute capability. The values that most affect the programming model are below, from Tables 30 and 31.

Note how much the data centre and consumer parts of one generation can differ. Compute capability 8.0 holds 2,048 threads per SM, while 8.6 holds 1,536. Occupancy tuned for one does not carry over to the other unchanged.

Compiling for compute capabilities

NVCC compiles device code to PTX, a virtual instruction set, and then to SASS, the machine code for a specific SM version. A binary can carry SASS for several real architectures plus PTX that the driver can compile just in time for GPUs that did not exist when you built it.

# SASS for A100 (8.0) and H100 (9.0), plus PTX for 9.0 so newer GPUs can JIT it
nvcc -O3 \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -gencode arch=compute_90,code=compute_90 \
  -o app app.cu

From compute capability 9.0, the guide describes three feature-set targets.

• The baseline target, such as compute_100, uses only features meant to carry forward and runs on 10.0 and later.
• The architecture-specific target, with an a suffix such as compute_90a or compute_100a, unlocks the full set of features specific to that one compute capability and runs only there.
• The family-specific target, with an f suffix such as compute_100f, unlocks the architecture-specific features shared by a family and runs on that family. The guide lists compute_100f as compatible with 10.0, 10.3 and 10.7, and compute_120f with 12.0 and 12.1.

The practical lesson: a kernel built with an a target is tied to one GPU generation. If you ship such a binary in a container, it will not run on the next one, and the failure appears at load time, not at build time.

What the current toolkit supports

The tables in the Programming Guide 13.4.2 cover compute capabilities from 7.5 upward. The CUDA Toolkit 13.4 Update 2 release notes record, for example, that the cuFFT library removed support for Maxwell, Pascal and Volta GPUs, "compute capabilities earlier than Turing". If you maintain code for older GPUs, pin the toolkit version that still supports them and record that pin, because it is part of your software inventory.

Hands-on 1: vector add, and why it is memory bound

Here is a complete vector add program in the style of the Programming Guide, with error checking and event timing. As noted at the top, this follows NVIDIA's documented API but was not compiled or run for this article.

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CHECK(call)                                                        \
    do {                                                                   \
        cudaError_t err_ = (call);                                         \
        if (err_ != cudaSuccess) {                                         \
            std::fprintf(stderr, "%s:%d %s\n", __FILE__, __LINE__,         \
                         cudaGetErrorString(err_));                        \
            std::exit(1);                                                  \
        }                                                                  \
    } while (0)

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

int main()
{
    const int n = 1 << 24;                     // 16,777,216 elements
    const size_t bytes = n * sizeof(float);    // 64 MiB per array

    float *hA = (float*)std::malloc(bytes), *hB = (float*)std::malloc(bytes),
          *hC = (float*)std::malloc(bytes);
    for (int i = 0; i < n; ++i) { hA[i] = 1.0f; hB[i] = 2.0f; }

    float *dA, *dB, *dC;
    CHECK(cudaMalloc(&dA, bytes));
    CHECK(cudaMalloc(&dB, bytes));
    CHECK(cudaMalloc(&dC, bytes));
    CHECK(cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice));

    const int block = 256;
    const int grid = (n + block - 1) / block;  // 65,536 blocks

    cudaEvent_t t0, t1;
    CHECK(cudaEventCreate(&t0));
    CHECK(cudaEventCreate(&t1));
    CHECK(cudaEventRecord(t0));
    vecAdd<<<grid, block>>>(dA, dB, dC, n);
    CHECK(cudaGetLastError());                 // catches bad launch configs
    CHECK(cudaEventRecord(t1));
    CHECK(cudaEventSynchronize(t1));

    float ms = 0.0f;
    CHECK(cudaEventElapsedTime(&ms, t0, t1));
    double gbps = 3.0 * bytes / (ms * 1e-3) / 1e9;   // 2 reads + 1 write
    std::printf("kernel %.3f ms, effective bandwidth %.1f GB/s\n", ms, gbps);

    CHECK(cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost));
    std::printf("C[0] = %f (expect 3.0)\n", hC[0]);

    cudaFree(dA); cudaFree(dB); cudaFree(dC);
    std::free(hA); std::free(hB); std::free(hC);
    return 0;
}

Build it with nvcc -O3 -arch=sm_80 vecadd.cu -o vecadd for an A100, or with the -gencode set above for several GPUs.

The roofline arithmetic

Each element does one floating point add and moves 12 bytes: two 4-byte reads and one 4-byte write. Its arithmetic intensity is therefore 1/12, about 0.083 FLOP per byte. Now compare that with the A100 figures from the Ampere whitepaper: 19.5 TFLOPS peak FP32 (non-tensor, at the 1,410 MHz boost clock) and 1,555 GB/s peak memory bandwidth. Both are peak figures, not sustained ones.

data moved            = 3 x 4 bytes x 16,777,216          = 201,326,592 bytes (201.3 MB)
lower bound on time   = 201,326,592 / 1,555e9 B/s          = 129.5 microseconds
ridge point           = 19.5e12 FLOP/s / 1,555e9 B/s       = 12.54 FLOP per byte
attainable at AI 1/12 = 1,555e9 x (1/12)                   = 129.6 GFLOP/s
share of FP32 peak    = 129.6e9 / 19.5e12                  = 0.66 %

However well it is written, the kernel cannot use more than about two thirds of one percent of the A100's FP32 peak; it would need about 150 times more arithmetic per byte to reach the ridge point. The right metric for this kernel is effective bandwidth, which the Best Practices Guide defines as bytes read plus bytes written, divided by time, and recommends as the primary measure for memory-bound code. The guide also warns to use the same divisor, 10^9 or 2^30, for theoretical and effective bandwidth, or the comparison is wrong.

Those are bounds computed from published peaks. We did not measure this kernel on an A100, and you should expect real effective bandwidth to land below the peak.

Two improvements are still worth knowing.

1. Grid-stride loop. Launch a grid sized to the GPU, a few blocks per SM, and let each thread loop with a stride of gridDim.x * blockDim.x. It removes the tail wave and makes the kernel correct for any n without changing the launch.
2. Vector loads. Reading float4 instead of float cuts the number of load instructions by four while keeping 100 percent sector efficiency, per the coalescing table. It needs 16-byte alignment and a remainder loop for sizes not divisible by four.

__global__ void vecAddStride(const float* A, const float* B, float* C, int n)
{
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;
         i += gridDim.x * blockDim.x)
        C[i] = A[i] + B[i];
}
// launch with, for example, grid = numSMs * 8, block = 256
// numSMs from cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount, dev)

Hands-on 2: reduction, and the steps that make it fast

Summing an array is the classic teaching kernel, because it touches every part of the model: coalescing, shared memory banks, divergence, synchronisation and the warp. The best known walkthrough is Mark Harris's NVIDIA presentation "Optimizing Parallel Reduction in CUDA". It reduces 4M (2^22) integers and reports times measured on a G80 GPU, whose peak bandwidth it gives as 86.4 GB/s from a 384-bit interface at 900 MHz DDR (384 times 1,800 divided by 8). Those are NVIDIA's measurements on hardware that is now about two decades old. We quote them for the relative effect of each step, not as numbers you would see today.

Source: Harris, "Optimizing Parallel Reduction in CUDA", NVIDIA Developer Technology, slide 34. We recomputed the bandwidth column from 2^22 times 4 bytes divided by each time and it matches to within rounding. Kernel 7 reaches 62.67 GB/s, or 72.5 percent of the G80's 86.4 GB/s peak. The same slide reports 73 GB/s for kernel 7 on 32M elements. Here is what each step changes.

Kernel 1: interleaved addressing with divergence

Each thread loads one element into shared memory. Then, in each of the log2(256) = 8 steps, thread tid adds its neighbour at distance s if tid % (2*s) == 0.

for (unsigned int s = 1; s < blockDim.x; s *= 2) {
    if (tid % (2 * s) == 0)
        sdata[tid] += sdata[tid + s];
    __syncthreads();
}

In the first step only even threads work, so every warp is divergent with half its lanes idle. In later steps, active threads are scattered across all warps, so every warp keeps issuing instructions for a handful of lanes. The slide also notes that the % operator is slow.

Kernel 2: strided index, no divergence, new bank conflicts

The fix is to make the active threads contiguous: thread tid works on index 2 * s * tid.

for (unsigned int s = 1; s < blockDim.x; s *= 2) {
    unsigned int index = 2 * s * tid;
    if (index < blockDim.x)
        sdata[index] += sdata[index + s];
    __syncthreads();
}

Now whole warps are active or idle, so divergence is gone, and the kernel runs 2.33 times faster. But look at the addresses: lane i touches word 2*s*i. That is a stride of 2s words, which by the bank table is a 2-way conflict at s = 1, 4-way at s = 2, and worse as s grows. We traded one problem for another.

Kernel 3: sequential addressing

Reverse the loop: start with a stride of half the block and halve it each step. Thread tid adds element tid + s.

for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (tid < s)
        sdata[tid] += sdata[tid + s];
    __syncthreads();
}

Lane i now reads words i and i + s, both stride 1 across the warp, so there are no bank conflicts, and active threads are contiguous, so there is no divergence beyond the last partial warp. Another 2.01 times. The cost is that half the threads are idle from the very first step.

Kernel 4: first add during the load

So halve the number of blocks and have each thread add two elements as it loads them from global memory. The global reads are still coalesced, the first reduction step is done in registers, and no thread is idle during the load. Another 1.78 times.

Kernel 5: unroll the last warp, and the trap

When s drops to 32 or below, only one warp is active. Harris's insight was that __syncthreads() and the loop overhead are wasted there, because, on the G80, the 32 threads of a warp executed in lockstep. His kernel unrolled the last six steps into straight-line code over a volatile shared memory pointer, with no synchronisation at all.

// From the 2007-era slides: correct under lockstep warps, NOT safe on CC 7.0+
__device__ void warpReduce(volatile int* sdata, int tid) {
    sdata[tid] += sdata[tid + 32];
    sdata[tid] += sdata[tid + 16];
    sdata[tid] += sdata[tid +  8];
    sdata[tid] += sdata[tid +  4];
    sdata[tid] += sdata[tid +  2];
    sdata[tid] += sdata[tid +  1];
}

This is exactly the "synchronisation-free intra-warp reduction" the current Programming Guide names as code to revisit. Under independent thread scheduling, lanes of a warp are no longer guaranteed to execute each line together, so one lane can read sdata[tid + 16] before the lane that owns it has finished its previous add. The volatile keyword stops the compiler from caching values in registers, but it does nothing about the order in which lanes run. On current GPUs, write the last warp with shuffles, or insert __syncwarp() between the steps.

// Safe on all current GPUs: warp-level reduction in registers
__device__ int warpReduceSum(int v) {
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xFFFFFFFF, v, offset);
    return v;
}

Kernels 6 and 7: unroll everything, then do more per thread

Kernel 6 makes the block size a template parameter, so the compiler knows it and can unroll the whole loop and drop dead branches. Kernel 7 has each thread sum many elements in a grid-stride loop before the block reduction begins. That is the same idea as the grid-stride vector add: fewer blocks, more work per thread, and the expensive tree reduction runs once per block instead of once per 512 elements. The slides argue that this is still work efficient, O(N) work overall, while hiding latency better.

A modern version

Putting it together with the tools CUDA has added since: a grid-stride loop for the global pass, warp shuffles for the in-warp steps, a tiny shared array for one value per warp, and one atomic per block. It needs no volatile tricks and is correct under independent thread scheduling. As with all code here, it follows the documented APIs but has not been run.

__inline__ __device__ float warpSum(float v) {
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xFFFFFFFF, v, offset);
    return v;
}

// blockDim.x must be a multiple of 32 and at most 1024
__global__ void reduceSum(const float* __restrict__ in, float* out, size_t n)
{
    __shared__ float warpTotals[32];          // one slot per warp (max 32 warps)

    float sum = 0.0f;
    for (size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x; i < n;
         i += (size_t)gridDim.x * blockDim.x)
        sum += in[i];                          // coalesced, many elements per thread

    sum = warpSum(sum);                        // 5 shuffle steps

    int lane = threadIdx.x % 32;
    int warp = threadIdx.x / 32;
    if (lane == 0) warpTotals[warp] = sum;
    __syncthreads();

    if (warp == 0) {
        int nWarps = blockDim.x / 32;
        sum = (lane < nWarps) ? warpTotals[lane] : 0.0f;
        sum = warpSum(sum);
        if (lane == 0) atomicAdd(out, sum);    // one atomic per block
    }
}

Because floating point addition is not associative and blocks finish in no fixed order, the last bits of the result can differ run to run. For bitwise reproducibility, write per-block partial sums and reduce them in a fixed order in a second pass.

For production work, NVIDIA's libraries provide tuned reductions. Writing one yourself is how you learn what the hardware does to your own kernels.

Occupancy, worked

The Programming Guide defines occupancy as the ratio of active warps on an SM to the maximum number of active warps the SM supports. Higher occupancy gives the warp schedulers more warps to choose from when one stalls on memory, which is how GPUs hide latency. It is a means, not an end, and later posts in this series cover when lower occupancy with more work per thread wins.

Four limits decide how many blocks fit on an SM: threads, block slots, registers and shared memory. The smallest wins.

The guide's own examples (compute capability 10.0)

The guide works through a device with 32 block slots, 2,048 threads, 65,536 registers and 233,472 bytes (228 KB) of shared memory per SM.

• testKernel<<<512, 768>>>(): 768 threads per block, so only 2 blocks fit in 2,048 threads. Occupancy is 1,536 / 2,048 = 75 percent.
• testKernel<<<512, 32>>>(): the 32 block slots run out first, so 32 blocks of 32 threads give 1,024 threads. Occupancy is 1,024 / 2,048 = 50 percent.
• A kernel using 100 KB of shared memory per block fits only 2 blocks, because a third would need 300 KB, more than the 228 KB available.

Register-limited examples on A100 (compute capability 8.0)

The guide's tables give compute capability 8.0 65,536 registers, 2,048 threads (64 warps), 32 block slots and up to 164 KB of shared memory per SM. Take blocks of 256 threads, which is 8 warps.

Register counts here are multiples of 8, so a warp's allocation (registers times 32) is already a multiple of 256. The Best Practices Guide notes that allocations are rounded up to the nearest 256 registers per warp, at least on compute capability 7.0. Its own example shows the effect: on 7.0, 37 registers per thread with 128-thread blocks gives 12 resident blocks and 75 percent occupancy. We reproduced that: 37 times 32 is 1,184, rounded up to 1,280 per warp, 5,120 per 4-warp block, and 65,536 divided by 5,120 is 12.8, so 12 blocks or 48 of 64 warps. For anything beyond back-of-envelope work, use the occupancy calculator in Nsight Compute or the cudaOccupancyMaxActiveBlocksPerMultiprocessor API, which know the exact rules for each architecture.

Common mistakes

1. Block sizes that are not a multiple of 32. A block of 100 threads runs 4 warps with 28 dead lanes, 21.9 percent of its capacity wasted for its whole life.
2. Forgetting the bounds check. The grid is rounded up to whole blocks, so the last block usually has threads past the end of the array. Without if (i < n) they read and write out of bounds, which may corrupt memory silently.
3. __syncthreads() inside a divergent branch. Every thread in the block must reach the same barrier. Hangs and wrong answers follow if they do not.
4. Warp-synchronous code without __syncwarp(). Code written for pre-Volta lockstep warps, especially unrolled last-warp reductions using volatile, can produce wrong results on compute capability 7.0 and later.
5. Array-of-structures layouts in hot kernels. Reading one field of a 32-byte struct gives 12.5 percent sector efficiency. Switch to structure of arrays for fields that are read together across threads.
6. Power-of-two strides in shared memory. Reading a column of a 32 by 32 tile is a 32-way bank conflict. Pad to 33 or use a swizzle.
7. Assuming local memory is fast. It is in device memory. Dynamically indexed private arrays and register spills land there. Check --resource-usage output for spill stores and loads.
8. Not checking launch errors. A kernel launch does not return an error code. Call cudaGetLastError() after the launch and check the result of the next synchronising call, or bad configurations fail silently.
9. Timing with the host clock without synchronising. A host timer around an asynchronous launch measures the launch, not the kernel. Use events or a profiler.
10. Using the legacy default stream for overlap. Work in the legacy default stream synchronises with other blocking streams, which serialises what you meant to overlap. Use non-blocking streams and pinned host memory.
11. Building only for your own GPU. A binary with SASS for sm_80 and no PTX will not run on a GPU it does not cover. A binary built with an a target runs on exactly one compute capability. Include PTX for forward compatibility unless you have a reason not to.
12. Relying on block execution order. Spin-waiting on another block's flag can deadlock. Use separate launches or a cooperative launch.

The GPU software supply chain behind every kernel

A CUDA kernel is a few dozen lines. The software that gets it onto a GPU is a stack of separately versioned components, each with its own release cycle and security bulletins.

Three facts from NVIDIA's own documentation make this an inventory problem rather than just a patching one.

• The CUDA runtime is statically linked by default. The NVCC documentation states that the static CUDA runtime library is the default on every platform except Windows, which uses a hybrid loader. A statically linked libcudart does not show up in the operating system's package database. The version lives inside your binary, and you only know it if your build records it.
• CUDA libraries ship as ordinary Python packages. PyPI hosts NVIDIA's runtime libraries as wheels, for example nvidia-cuda-runtime (version 13.4.92 at the time of writing) and the CUDA 12 line as nvidia-cuda-runtime-cu12. A pip install of a framework can pull a whole CUDA library set as transitive dependencies, which then live in the Python dependency tree rather than in a system package list.
• Bulletins are moving to machine-readable formats. NVIDIA's Product Security page says that from 1 October 2026 its PSIRT publishes security bulletins on GitHub as Markdown, CSAF and CVE records, in parallel with the website. CSAF is the format automated vulnerability tooling can consume directly, which only helps if you know which NVIDIA components and versions you run.

The cost of all this is real. Our estimate of GPU cost per million tokens when self-hosting shows how much hardware sits behind a serving stack, and our walkthrough of how a GPU is made, from wafer to HBM package shows how much engineering is in the silicon. The software stack on top is where most of the day-to-day security work is.

Where Safeguard fits

Safeguard does not design, program or tune GPUs, and nothing in this article is something Safeguard makes faster. Kernel performance is between you, your compiler and your profiler. The fit is one layer out, in the GPU software supply chain described above.

Safeguard's software composition analysis builds a transitive inventory of dependencies across ecosystems including PyPI, and includes container image layers, and Safeguard generates SBOMs in SPDX and CycloneDX. That covers the parts of a GPU stack that arrive as Python wheels and as layers in a container image: the CUDA runtime and library wheels a framework pulls in, and the libraries baked into a CUDA base image. Vulnerability matching draws on sources including NVD, GHSA, OSV, CISA KEV and vendor bulletins, and reachability analysis helps rank which findings sit on code paths your application actually calls. What Safeguard does not see is a runtime statically linked into a binary that your build did not record, or a host kernel-mode driver outside the images you scan. Those need your build to emit the version and your host inventory to track the driver.

Each rebuild of a GPU image pulls the dependency tree again. Keeping fixed versions flowing into those rebuilds, and known-vulnerable ones out, is the same discipline as for any application, so the section below applies here too.

FAQ

What is the difference between a thread, a warp and a block in CUDA?

A thread is the unit you write code for. A warp is 32 threads that the hardware schedules together and that issue one instruction at a time. A block is a group of up to 1,024 threads, divided into warps, that runs on one SM and can share shared memory and synchronise with __syncthreads().

What is SIMT, and how is it different from SIMD?

SIMT, Single Instruction Multiple Thread, has you write scalar code for one thread while hardware runs 32 threads per instruction. SIMD exposes the vector width to software. In SIMT each thread may branch independently, and the hardware masks lanes to handle divergence.

What is warp divergence and how much does it cost?

When lanes of one warp take different branch paths, the warp runs each path with the other lanes masked. A two-way split runs both paths, so lane utilisation for that region falls to at most 50 percent. Divergence between different warps costs nothing extra.

What changed with independent thread scheduling in Volta?

From compute capability 7.0, each thread has its own program counter and call stack, and threads of a warp can diverge and reconverge at sub-warp granularity. Starvation-free algorithms such as fine-grained locks become possible within a warp. Code that assumed lockstep warps must use __syncwarp() or _sync intrinsics.

Is CUDA local memory fast?

No. Local memory is private to a thread in scope but lives in device memory, with global memory latency and bandwidth. Register spills and dynamically indexed private arrays go there. It is cached and is coalesced when all lanes access the same relative address.

What is memory coalescing in CUDA?

Global memory is served in 32-byte sectors. Coalescing is the hardware combining a warp's 32 addresses into as few sectors as possible. Consecutive 4-byte accesses need 4 sectors, 100 percent efficient. Accesses 32 bytes or more apart need 32 sectors, 12.5 percent efficient.

How do I avoid shared memory bank conflicts?

Shared memory has 32 four-byte banks. Lane strides with a large common factor with 32 conflict. Use stride 1 or any odd stride, pad 2D tiles to an odd row width such as 33, or use an XOR swizzle on the column index.

What are thread block clusters?

A level of hierarchy added in compute capability 9.0. Blocks in a cluster are co-scheduled on one GPC and can access each other's shared memory, called distributed shared memory. The portable maximum is 8 blocks, and H100 allows 16 with an opt-in attribute.

More in this series

Other posts in our GPUs, CPUs, QPUs and lithography series:

• AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP
• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL
• CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared
• GPU Software Supply Chain: Drivers, CUDA, cuDNN, NCCL and Your SBOM
• How a QPU Works: Qubit Chips, Fridges, Control Pulses and Readout
• How CPUs Work: Pipelines, Out-of-Order, Caches, Branches, SIMD
• How EUV Lithography Machines Work: Tin Plasma, Mirrors and High-NA
• How GPUs Work: Graphics Pipeline, Shader Cores, Ray Tracing and Compute
• Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores
• Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL
• NVIDIA GPU Architecture History: From Tesla and Fermi to Blackwell
• NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared
• Optimizing a CUDA Matrix Multiply Step by Step, From Naive to Tensor Cores
• Photolithography Explained: Contact Printing to DUV Immersion and Overlay
• Tile-Based GPUs Explained: Apple, Mali, Adreno, PowerVR and TBDR

Fix it without waiting for upstream

Whatever the vulnerability, most of the risk sits in two gaps. The first is the time between a fix being published and that fix actually running in your production builds. The second, and worse, is the time before any fix exists at all, while a maintainer works out a patch and a release schedule. Waiting through both gaps is a habit, not a requirement. Safeguard gives teams three ways to close them.

1. Autonomous remediation: pull requests, not tickets

When a fixed version exists, Safeguard's Griffin models open a tested, policy gated pull request that moves the affected dependency to the safe version. You can run it on one repository or in bulk across every project that carries the vulnerable version. Remediation covers eight categories, including dependencies, licenses, security posture and findings. Reachability analysis ranks the vulnerable code paths your application actually calls first, so the first pull requests you merge are the ones that matter most. Autonomous remediation ships switched off, with every category set to Manual, so your team decides how much runs on its own and what still needs a human approval.

2. Gold Open Source: the safe version in seconds, free

gold.safeguard.sh is free and needs no login. Search any CVE or package to see affected and fixed versions, CVSS, EPSS, CISA KEV status and a recommended safe version, or use the "Am I affected?" checker. The same data is available through a free read API, README status badges and RSS threat feeds. The safeguard-sh/gold-scan-action GitHub Action can fail a build on a known exploited vulnerability, on any known vulnerability, or never, so a version you have just removed cannot quietly come back in next week's merge.

3. Gold Registry: start from components that are already fixed

registry.safeguard.sh provides hardened, drop in replacements for upstream packages and container images. Every Gold artifact is held continuously to zero exploitable critical or high CVEs, zero malware and a completed license check, with SLSA L3 or higher provenance, a Sigstore signature and a public evidence page you can hand to an auditor. An artifact loses its Gold badge if it regresses. Teams can point their own package directory at the registry so hardened equivalents resolve without changing how developers install anything.

When upstream has no fix yet, or will not ship one, Safeguard first works to get the fix accepted upstream. If the issue is severe or upstream will not move, Safeguard builds a fork, fixes the vulnerability there and publishes a hardened Gold build, so you are not stuck on a maintainer's release schedule. You can also request a hardened, compatibility tested build of a specific package. This fork and patch flow is in early access today, so ask which packages it covers before you plan around it.

One honest limit: when the vulnerable code lives in a vendor appliance, a hosted service or an operating system runtime you do not build yourself, the fix still has to come from that vendor or distribution. The steps above apply to the open source dependencies and container images you build and ship, which is where most teams carry the largest and least tracked share of their exposure.

See how Safeguard would help

The first demo is a real Safeguard CLI session. You install it, run safeguard to open the interactive terminal and ask it to fix the lodash CVE in a repository. It reads the finding and advisory, asks before any command that changes something, edits package.json to lodash 4.17.21, runs npm install and npm test with your approval, prints a summary and pushes nothing. The demo uses an npm project, and the same flow applies to the Python requirements of a GPU training or serving repository; /scan . gives you a first inventory of what it pulls in.

Click through it at your own pace, or open it full screen.

The second demo sets up the local runner on your own machine, approves it with its device code, then scans from the terminal. GPU build hosts often sit on private networks next to large model weights and datasets, and the runner lets scans happen where the code and images already are.

Click through it at your own pace, or open it full screen.

The third demo connects Claude to the Safeguard MCP server with OAuth and shows how you choose which tools Claude can use. Once connected, you can ask which of your projects contain a given component and version, which is the question to answer when a new bulletin names a CUDA library.

Click through it at your own pace, or open it full screen.

To try it on your own projects, download Safeguard or read the CLI documentation.

Quick check

Sources

• NVIDIA CUDA Programming Guide 13.4.2: Programming Model
• NVIDIA CUDA Programming Guide: Writing SIMT Kernels
• NVIDIA CUDA Programming Guide: Unified and System Memory
• NVIDIA CUDA Programming Guide: Asynchronous Execution
• NVIDIA CUDA Programming Guide: Advanced Kernel Programming (SIMT, independent thread scheduling)
• NVIDIA CUDA Programming Guide: Cooperative Groups
• NVIDIA CUDA Programming Guide: Asynchronous Data Copies
• NVIDIA CUDA Programming Guide: Unified Memory
• NVIDIA CUDA Programming Guide: C++ Language Extensions
• NVIDIA CUDA Programming Guide: Compute Capabilities
• NVIDIA CUDA C++ Best Practices Guide
• NVIDIA Hopper Tuning Guide
• NVIDIA CUDA Compiler Driver NVCC documentation
• NVIDIA CUDA Toolkit 13.4 Update 2 Release Notes
• NVIDIA PTX ISA
• NVIDIA Tesla V100 GPU Architecture whitepaper (WP-08608-001 v1.1, August 2017)
• NVIDIA A100 Tensor Core GPU Architecture whitepaper (2020)
• Mark Harris, "Optimizing Parallel Reduction in CUDA", NVIDIA Developer Technology
• NVIDIA Product Security and Security Bulletins
• PyPI: nvidia-cuda-runtime
• PyPI: nvidia-cuda-runtime-cu12

Last reviewed: 25 October 2026. CUDA documentation is updated with each toolkit release; check the linked NVIDIA pages for the version you build with before relying on a specific limit.