Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores

An NVIDIA streaming multiprocessor (SM) is the unit that actually runs your CUDA threads. Since Volta, each SM is split into four partitions, and each partition has one warp scheduler, a 64 KB slice of the register file, its own math units and, on recent parts, one tensor core. The four partitions share an L1 cache that doubles as shared memory. Occupancy is simply how many warps the SM holds at once, out of a hardware maximum that is 64 warps on A100 and H100.

Everything else in GPU performance tuning follows from those few facts. Registers and shared memory are a fixed budget per SM, so the more each thread uses, the fewer warps fit. Fewer warps means fewer candidates for the scheduler when one warp is waiting on memory. Whether that matters depends on how much independent work each warp carries, which is why a matrix multiply kernel can run at 12.5 percent occupancy and still be close to peak while a simple copy kernel at the same occupancy crawls.

Our sources are NVIDIA's architecture whitepapers, the CUDA Programming Guide (version 13.4.2), the Best Practices Guide, the tuning guides, the PTX ISA (version 9.4) and the Nsight Compute Profiling Guide (version 2026.3.1). Every computed number was checked with Python. We could not run CUDA code on the Mac we wrote this on, so the code samples follow NVIDIA's documentation but were not executed, and we report no benchmarks of our own.

Key takeaways

• Since Volta, an SM is four partitions, each with one warp scheduler, one dispatch unit, an L0 instruction cache and a 64 KB register file.
• The register file is 64K 32-bit registers (256 KB) per SM on every compute capability from 7.5 to 12.x, with at most 255 registers per thread.
• A100 and H100 both allow 64 resident warps and 32 blocks per SM. Shared memory per SM is up to 164 KB on A100 and 228 KB on H100, minus 1 KB that CUDA reserves per block.
• Registers are allocated in units of 256 per warp, so 33 registers per thread costs the same as 40. At 1,024 threads per block, that one extra register halves occupancy.
• By Little's law, an H100 SM needs roughly 6.9 KB of loads in flight to use its share of HBM bandwidth, assuming a 500 cycle memory latency.
• High occupancy is a means, not a goal. Kernels with plenty of independent work per warp can saturate the hardware with few warps.
• In Nsight Compute, read theoretical occupancy, achieved occupancy, eligible warps per scheduler and stall reasons, in that order.

What a streaming multiprocessor is, and where it sits

A CUDA GPU is a hierarchy. At the top level, the chip has a number of graphics processing clusters (GPCs). Each GPC holds texture processing clusters (TPCs), and each TPC holds two SMs. The A100 whitepaper describes the full GA100 die as having 128 SMs, of which the shipping A100 product enables 108, arranged as 7 GPCs with 7 or 8 TPCs each. The H100 whitepaper describes a full GH100 die with 144 SMs, of which the H100 SXM5 product enables 132 and the PCIe card enables 114.

GPU
├── L2 cache (A100: 40 MB, H100 SXM5: 50 MB)
├── Memory controllers to HBM (A100: 5 stacks, H100 SXM5: HBM3)
└── GPCs
    └── TPCs (2 SMs each)
        └── SM
            ├── 4 x partition (processing block, "SMSP")
            │   ├── L0 instruction cache
            │   ├── warp scheduler + dispatch unit
            │   ├── register file, 64 KB (16,384 x 32-bit)
            │   ├── INT32, FP32, FP64 units
            │   ├── tensor core
            │   └── load/store and special function units
            └── shared by all 4 partitions
                ├── L1 data cache / shared memory (A100: 192 KB, H100: 256 KB)
                ├── texture units
                └── RT core (graphics parts only)

The SM is where scheduling decisions happen. When you launch a kernel with a grid of thread blocks, the hardware block scheduler places whole blocks onto SMs. A block never splits across SMs. Inside the SM, each block is cut into warps of 32 threads, and the Nsight Compute Profiling Guide states that a warp is allocated to one sub-partition and stays there "from launch to completion". From then on, the warp competes with other warps on the same partition for the attention of that partition's scheduler.

The SM count sets how many blocks run at once. On A100 at 8 blocks per SM, one wave is 864 blocks, so a grid of 900 blocks needs a second wave of only 36 blocks, and most of the GPU idles while it runs.

The SM, one partition at a time

NVIDIA introduced the four partition design with Volta. The Volta whitepaper contrasts it with Pascal's GP100 SM, which had two processing blocks with one warp scheduler, two dispatch units and a 128 KB register file each. The GV100 SM has four, each with 16 FP32 cores, 8 FP64 cores, 16 INT32 cores, two tensor cores, a new L0 instruction cache, one warp scheduler, one dispatch unit and a 64 KB register file.

That shape has held since. The Turing and GA102 whitepapers describe the same four blocks, each with a 64 KB register file, an L0 instruction cache, one warp scheduler and one dispatch unit. The Nsight Compute guide calls each one an SM sub-partition (SMSP) and lists what the four share: the unified L1 data cache and shared memory, the texture units, and RT cores where they exist.

Here is one partition of an A100 SM drawn from the whitepaper's figures and text. Unit counts come from the whitepaper's per-SM totals divided by four, which is how NVIDIA draws them.

          One A100 SM partition (one quarter of the SM)
+------------------------------------------------------------+
| L0 instruction cache                                       |
| Warp scheduler (1 warp per clock) | Dispatch unit          |
+------------------------------------------------------------+
| Register file: 16,384 x 32-bit = 64 KB                     |
+------------------------------------------------------------+
| 16 x INT32 | 16 x FP32 | 8 x FP64 | 1 x 3rd gen tensor core |
+------------------------------------------------------------+
| Load/store units           | Special function unit (SFU)   |
+------------------------------------------------------------+
            |                                  |
            v                                  v
+------------------------------------------------------------+
| Shared by all four partitions:                             |
| 192 KB L1 data cache / shared memory, texture units        |
+------------------------------------------------------------+

The whitepapers draw load/store units and special function units inside every partition, but the text of the A100 and H100 papers does not state how many of each there are, so we leave the counts out rather than read them off a picture.

The warp scheduler and dispatch unit

Each partition's scheduler keeps a pool of resident warps. The Nsight Compute guide's Scheduler Statistics section describes the loop precisely: on every cycle, each scheduler checks the state of the warps allocated to it. Active warps that are not stalled are called eligible. The scheduler selects one eligible warp and issues one or more instructions from it. If no warp is eligible, the issue slot is skipped and nothing is issued that cycle.

The pool has a fixed size. The Nsight Compute guide gives 8 warps for Turing, which matches the 32 resident warps per SM that the Programming Guide allows at compute capability 7.5, split four ways. By the same arithmetic, A100 and H100 have 16 warp slots per scheduler and 8.6 or 8.9 parts have 12.

For a warp to be eligible, the Nsight Compute guide lists three conditions: the warp has a decoded instruction, all of that instruction's input dependencies are resolved, and the function unit it needs is available. Everything that stops a warp from meeting those conditions shows up in the profiler as a stall reason, which we cover later.

Since Volta, threads in a warp need not run in lockstep. The CUDA Programming Guide explains that independent thread scheduling keeps a program counter and call stack per thread. Divergent branches still cost throughput, and code that assumed implicit warp synchronous behavior needs explicit __syncwarp() calls.

The register file

Registers are the fastest storage in the SM and the most contested. The CUDA Programming Guide's memory table gives 64K 32-bit registers per SM for every compute capability from 7.5 through 12.x, a maximum of 64K registers per thread block, and a maximum of 255 registers per thread. 65,536 registers of 4 bytes each is 256 KB, which is the "Register File Size / SM" line in both the A100 and H100 whitepaper tables. Split four ways, each partition has 16,384 registers, the 64 KB per partition the whitepapers describe.

Across the GPU, that is 27,648 KB on A100 (108 SMs) and 33,792 KB on H100 SXM5 (132 SMs), as both whitepapers list. The register file is larger than the shared memory on both chips.

The Best Practices Guide says register access "consumes zero extra clock cycles per instruction", apart from read after write dependencies and bank conflicts. The catch is capacity: for 2,048 resident threads, each can have at most 65,536 divided by 2,048, or 32 registers.

Two details make the arithmetic less obvious than that division. First, the Best Practices Guide says "register allocations are rounded up to the nearest 256 registers per warp". A warp of 32 threads at 33 registers needs 1,056 registers, which rounds up to 1,280, the same as 40 registers per thread. Second, the register file is physically four 64 KB banks, one per partition, and each warp lives on one partition. We show below that treating registers as a per-partition budget reproduces the Best Practices Guide's own worked example exactly.

When a kernel's live values do not fit in its allowed registers, the compiler spills them to local memory, which lives in device memory and is cached in L1 and L2. The Programming Guide notes that capping registers with -maxrregcount can let more blocks fit but "may also result in more register spilling", and that the trade is sometimes still a net win.

L1 data cache, shared memory and the carveout

Since Volta, the L1 data cache and shared memory are one SRAM array, and you choose how much of it acts as shared memory. The Ampere tuning guide states that A100 raised the combined L1, texture cache and shared memory capacity to 192 KB per SM, 50 percent more than V100, and that shared memory per SM can be 0, 8, 16, 32, 64, 100, 132 or 164 KB. The Hopper tuning guide raises the combined capacity to 256 KB per SM and the shared memory maximum to 228 KB, a 39 percent increase over A100's 164 KB. Compute capability 8.6 and 8.9 parts, the graphics oriented Ampere and Ada chips, have a 128 KB combined array and up to 100 KB of shared memory.

Three rules from the tuning guides matter for occupancy:

1. CUDA reserves 1 KB of shared memory per thread block. That is why the per-block maximum is 163 KB on A100 and 227 KB on H100, one less than the per SM figure.
2. Static shared memory declarations are limited to 48 KB. Anything larger has to be dynamic shared memory with an explicit opt-in through cudaFuncSetAttribute.
3. The split between L1 and shared memory is a preference you set with cudaFuncAttributePreferredSharedMemoryCarveout. The driver picks one of the supported capacities.

The shared memory array has 32 banks on every compute capability in the current table. When the 32 threads of a warp access 32 different banks, the access completes in one pass. When several threads hit different addresses in the same bank, the accesses serialize. This bank structure is why so many CUDA kernels pad their shared memory tiles by one column.

Hopper added a way to reach other SMs' shared memory. The H100 whitepaper introduces thread block clusters, groups of blocks guaranteed to run concurrently on SMs within one GPC, and distributed shared memory, which "allows direct SM-to-SM communications for loads, stores, and atomics". The Blackwell tuning guide gives a portable maximum cluster size of 8, and 16 on B200 with an opt-in.

Load/store units, the special function unit and the texture path

Every memory instruction from a warp goes through the partition's load/store path into the shared L1. The L1 acts as what the tuning guides call a coalescing buffer: it gathers the addresses requested by the 32 threads of a warp before delivering the data. A warp that reads 32 consecutive 4-byte floats needs 128 bytes from a small number of cache lines. A warp that reads 32 floats spaced 4 KB apart touches 32 different lines.

Ampere added an asynchronous copy from global memory into shared memory that, per the A100 whitepaper, skips the trip through the register file and overlaps with computation. Hopper's Tensor Memory Accelerator (TMA) goes further: the H100 whitepaper describes it as a unit that "can efficiently transfer large blocks of data between global memory and shared memory", including between blocks in a cluster.

The special function unit runs special math instructions. The Nsight Compute guide groups special math, dynamic branches and shared memory instructions into the MIO pipelines, so a kernel heavy on fast transcendentals can show the same throttle signature as one that hammers shared memory.

Tensor cores

Tensor cores are matrix multiply and accumulate units. Each takes small matrix tiles and computes D equals A times B plus C in a few instructions instead of thousands of scalar FMAs. The A100 whitepaper says Volta and Turing have eight tensor cores per SM, each doing 64 FP16/FP32 FMA operations per clock, while A100 has four per SM, each doing 256, for 1,024 dense FMAs per SM per clock. The H100 whitepaper says Hopper's fourth generation tensor core doubles A100's per SM rate clock for clock on the same data types, and quadruples it with FP8.

We checked those per SM claims against the peak throughput tables. A100's dense FP16 tensor peak is 312 TFLOPS at a 1,410 MHz boost clock across 108 SMs. Dividing 312 trillion by 2 (an FMA counts as two operations), by 108 and by 1.41 billion gives 1,024 FMAs per SM per clock. H100 SXM5's dense FP16 tensor peak is 989.4 TFLOPS at the 1,830 MHz boost clock the whitepaper lists for tensor operations across 132 SMs, which gives 2,048 FMAs per SM per clock, or 512 per tensor core. Its FP8 peak of 1,978.9 TFLOPS gives 4,096.

The programming interface has moved with the hardware. Volta exposed tensor cores through warp level matrix functions in CUDA C++ (wmma), where one warp cooperatively owns each tile. The mma.sync PTX instructions used from Volta through Ampere are also warp level. Hopper added wgmma.mma_async, which the PTX ISA defines over a warpgroup, "a set of four contiguous warps such that the warp-rank of the first warp is a multiple of 4." One warpgroup naturally spans all four partitions of an SM, so all four tensor cores work on one large tile together, and the instruction is asynchronous so the warps can issue loads while the math runs.

The fifth generation tensor core in compute capability 10.0 parts goes further. The PTX ISA says it "has dedicated on-chip memory that is specialized for use by TensorCore operations", called Tensor Memory, and that on sm_100a and sm_100f it is organized as 512 columns by 128 rows per CTA of 32-bit cells. That works out to 256 KiB. Tensor Memory is allocated by one warp in units of 32 columns and must be freed before the kernel exits. The Nsight Compute guide adds that on 5th generation MMA, one issuing thread can drive work across all four warp schedulers. The upshot for this post is that on Blackwell data center parts, the accumulator of a big matrix multiply no longer has to live in the general register file, which changes the register arithmetic for those kernels.

RT cores on graphics parts

A100 has no RT cores; its whitepaper says so. Graphics parts do. The Turing whitepaper introduced one RT core per SM to accelerate bounding volume hierarchy (BVH) traversal and ray and triangle intersection, and the GA102 whitepaper's second generation RT core can run concurrently with graphics or compute work on the same SM. For compute kernels that never trace rays, the RT core sits idle and does not affect occupancy.

SM anatomy per generation

The table below collects the per SM figures that the whitepapers and the CUDA Programming Guide state directly. Where a source gives only a GPU total, we divided by the SM count and say so.

The V100 and A100 columns come from the A100 whitepaper's comparison table; the V100 warp limit comes from the Best Practices Guide, since the current Programming Guide table starts at 7.5. TU102 figures come from the Turing whitepaper, which shows 72 SMs on the full die. We leave out the GA102 tensor rate because it differs between consumer and workstation boards. The Blackwell column uses only the Blackwell tuning guide and the Programming Guide, which do not list per SM lane counts, so we do not guess them.

One more line from the tables is worth knowing. The CUDA Programming Guide's table gives a 2:1 FP32 to FP64 throughput ratio for compute capability 8.0 and 9.0, and 64:1 for the graphics oriented 8.6 and 8.9 parts. A kernel that does serious double precision work will behave completely differently on those two families even though they share the Ampere name.

How a warp moves through the SM

A warp's life on one partition can be described as a loop through a few states. The Nsight Compute guide says a warp is stalled when it waits on an instruction fetch, a memory dependency, an execution dependency or a synchronization barrier. Otherwise it is eligible, and if the scheduler picks it, it is selected and issues.

Picture one scheduler with four warps, W0 to W3, running a loop that loads a value from global memory, does three arithmetic instructions on it, and stores a result. The timeline below is illustrative, not measured. It assumes, as the Best Practices Guide states for compute capability 7.0, that most arithmetic instructions have about 4 cycles of latency, and it uses a placeholder of 20 cycles for the memory load so the picture fits on screen. Real global memory latency is, in the Best Practices Guide's words, "hundreds of clock cycles".

cycle:  0    1    2    3    4    5    6    7    8   ...  20   21   22   23   24
W0:    LD   ~~~~~~~~~~~~~~~~ waiting on memory ~~~~~~~~~~~~  FMA  .    .    .   FMA
W1:    .    LD   ~~~~~~~~~~~~~~~~ waiting on memory ~~~~~~~~~~~~  FMA  .    .
W2:    .    .    LD   ~~~~~~~~~~~~~~~~ waiting on memory ~~~~~~~~~~~~  FMA  .
W3:    .    .    .    LD   ~~~~~~~~~~~~~~~~ waiting on memory ~~~~~~~~~~~~  FMA
issue: W0   W1   W2   W3   --   --   --   --   --  ...  W0   W1   W2   W3   W0

"--" = issue slot skipped: no eligible warp

With four warps and a 20 cycle load, the scheduler issues four loads and then idles for 16 cycles. Each skipped slot is wasted issue bandwidth. If there were 20 warps, each with its own independent load, the scheduler would have something to issue every cycle. That is latency hiding by warp switching: the SM does not make memory faster, it overlaps the waits.

Switching costs nothing, because each resident warp's registers stay in the register file for its whole life. There is no save and restore as in a CPU context switch, which is exactly why registers are so precious: a waiting warp still holds its full allocation.

Latency hiding with numbers: Little's law

Little's law says the average number of items in a system equals the arrival rate times the average time each spends in it. For a GPU, the items are bytes or instructions in flight. To keep a pipe with throughput T busy when each item takes latency L, you need T times L items in flight.

Arithmetic latency

Take the Best Practices Guide's figure of roughly 4 cycles of latency for most arithmetic instructions on compute capability 7.0. Each partition's scheduler can issue for one warp per clock. To issue every cycle with a 4 cycle dependency, the scheduler needs about 4 independent instructions in flight. Those can come from 4 different warps, or from fewer warps that each have independent instructions back to back. The Guide's conclusion is that this latency "can be completely hidden by the execution of threads in other warps." On an SM with four schedulers, that means at least 16 warps in total if each warp has no independent instructions of its own.

Memory latency

Memory is the harder case because the latency is much longer and the throughput target is per byte. NVIDIA's guides do not publish a single global memory latency figure in cycles; they say "hundreds". So we pick an assumption, 500 cycles, and label it clearly as one.

H100 SXM5's peak memory bandwidth in the whitepaper's table is 3,352 GB/s. Divide by 132 SMs and by the 1.83 GHz boost clock and each SM's share is about 13.9 bytes per SM clock. With an assumed 500 cycle latency, each SM needs about 6,940 bytes of loads outstanding at all times. A warp that loads one 4-byte float per thread moves 128 bytes, so that is about 54 warps each with one load in flight. If each thread loads a float4 instead, a warp moves 512 bytes per instruction and the requirement drops to about 14 warps.

Two conclusions follow. A kernel doing one scalar load per thread per loop needs most of the 64 warp slots just to cover memory latency, so occupancy matters for it. And the same bandwidth can be reached with a quarter of the warps if each thread has more bytes in flight. Because the bandwidths are vendor peaks and the latency is our assumption, treat these warp counts as orders of magnitude.

Occupancy from first principles

The CUDA Programming Guide defines occupancy as "the ratio of the number of active warps to the maximum number of active warps supported by the SM." The Nsight Compute guide uses the same definition and adds the important caveat: "Higher occupancy does not always result in higher performance, however, low occupancy always reduces the ability to hide latencies."

There are two kinds:

• Theoretical occupancy is what the launch configuration allows. It depends only on the block size, the registers per thread, the shared memory per block and the hardware limits. You can compute it before running anything.
• Achieved occupancy is the average number of active warps actually measured during the run, divided by the maximum. It is lower than theoretical when the grid is too small to fill every SM, when blocks finish at very different times, or in the tail of the last wave.

Theoretical occupancy is the minimum of four limits:

1. Warps. The maximum resident warps per SM divided by warps per block.
2. Blocks. The maximum resident blocks per SM.
3. Registers. How many blocks fit in the register file, after rounding.
4. Shared memory. How many blocks fit in the shared memory carveout, counting the 1 KB per block that CUDA reserves.

Hopper and later add a fifth, for kernels launched in clusters: the Blackwell tuning guide recommends computing occupancy with cudaOccupancyMaxActiveClusters for those.

The hardware limits for the generations people tune most, taken from the CUDA Programming Guide's compute capability tables, are:

One discrepancy to flag. The CUDA Programming Guide's table lists 24 resident blocks per SM for compute capability 12.x, while the Blackwell tuning guide says the maximum is 32 for both 10.0 and 12.0. The tuning guide also says shared memory capacity per SM is 128 KB for 12.0, while the Programming Guide lists 100 KB of maximum shared memory per SM, with 128 KB as the size of the unified data cache. We use the Programming Guide's figures for cc 12.x and recommend checking cudaGetDeviceProperties on the actual device.

Occupancy maths, worked for A100 and H100

The method

Here is the calculation as a short Python function. It follows the documented rules: registers rounded up to 256 per warp, 1 KB of shared memory reserved per block, and the hardware limits from the table above. For the register limit, it treats the register file as four per-partition banks of 16,384 registers, each holding whole warps. That is how the hardware is built according to the whitepapers, and it is the only reading we found that reproduces the Best Practices Guide's published example, as we show next.

import math

LIMITS = {
    "A100 (cc 8.0)": dict(warps=64, blocks=32, regs=65536, smem=164 * 1024),
    "H100 (cc 9.0)": dict(warps=64, blocks=32, regs=65536, smem=228 * 1024),
    "cc 8.6":        dict(warps=48, blocks=16, regs=65536, smem=100 * 1024),
    "cc 7.0":        dict(warps=64, blocks=32, regs=65536, smem=96 * 1024),
}

def occupancy(gpu, threads_per_block, regs_per_thread, smem_per_block):
    L = LIMITS[gpu]
    warps_per_block = math.ceil(threads_per_block / 32)

    # 1. warp slots
    by_warps = L["warps"] // warps_per_block
    # 2. block slots
    by_blocks = L["blocks"]
    # 3. registers: rounded up to 256 per warp, four partition banks
    regs_per_warp = math.ceil(regs_per_thread * 32 / 256) * 256
    warps_by_regs = (L["regs"] // 4 // regs_per_warp) * 4
    by_regs = warps_by_regs // warps_per_block
    # 4. shared memory: 1 KB reserved per block
    by_smem = L["smem"] // (smem_per_block + 1024) if smem_per_block else 10**9

    blocks = min(by_warps, by_blocks, by_regs, by_smem)
    warps = blocks * warps_per_block
    return blocks, warps, 100 * warps / L["warps"]

This is a teaching model, not a replacement for the real calculator. Nsight Compute's occupancy section and the CUDA runtime's occupancy API know the exact allocation granularities for each architecture, including the shared memory allocation unit and the carveout the driver will actually choose. Use them for decisions.

Checking the model against NVIDIA's own example

The Best Practices Guide gives this example for compute capability 7.0: "a kernel with 128-thread blocks using 37 registers per thread results in an occupancy of 75% with 12 active 128-thread blocks per multi-processor, whereas a kernel with 320-thread blocks using the same 37 registers per thread results in an occupancy of 63% because only four 320-thread blocks can reside on a multiprocessor."

Walk through it. 37 registers times 32 threads is 1,184 registers per warp, which rounds up to 1,280. Each partition has 16,384 registers, and 16,384 divided by 1,280 is 12.8, so 12 warps fit per partition and 48 per SM.

• With 128-thread blocks, each block is 4 warps. 48 divided by 4 is 12 blocks, 48 warps, and 48 of 64 is 75 percent. That matches.
• With 320-thread blocks, each block is 10 warps. 48 divided by 10 is 4 whole blocks, 40 warps, and 40 of 64 is 62.5 percent, which the Guide rounds to 63. That matches too.

If you instead divide the whole 65,536 registers by 1,280, you get 51 warps, which would allow 5 blocks of 320 threads and 78 percent occupancy. The Guide says 4. The per-partition reading is the one that agrees with NVIDIA's published numbers.

Worked cases on A100 and H100

We ran the function on ten configurations for A100 and H100. These are theoretical occupancy figures from the model, not measurements.

A few rows deserve a closer look.

The 33 register cliff. At 1,024 threads per block and 32 registers, a warp needs exactly 1,024 registers, each partition holds 16 warps, the SM holds 64, and two 32-warp blocks fit. Add one register per thread and each warp rounds up to 1,280 registers. Each partition now holds 12 warps, the SM holds 48, and only one 32-warp block fits, because a block cannot be split across the gap. Occupancy falls from 100 percent to 50 percent on both A100 and H100 for the cost of one register. This is the kind of drop that appears after an innocent code change and is invisible unless you read the compiler's register report.

Shared memory heavy kernels. A block of 128 threads that uses 48 KB of shared memory actually costs 49 KB once the 1 KB reservation is added. 164 KB divided by 49 KB is 3 blocks on A100, and 228 KB divided by 49 KB is 4 blocks on H100. The bigger Hopper carveout buys a third more resident blocks for the same code.

When the limiter changes between GPUs. At 128 threads, 96 registers and 32 KB of shared memory, A100 runs out of shared memory first, at 4 blocks of 33 KB, while H100 has shared memory for 6 blocks but runs out of registers at 5. Each warp needs 3,072 registers, 5 warps fit per partition, 20 per SM, and 20 divided by 4 warps per block is 5 blocks. Raising the H100 shared memory allocation will not help that kernel; reducing registers would.

Small blocks. 64-thread blocks of 2 warps reach 100 percent on A100 and H100 only because those parts allow 32 resident blocks. The same configuration on a compute capability 8.6 part, which allows 16 blocks and 48 warps, stops at 16 blocks, 32 warps and 66.7 percent. Tiny blocks are an easy way to hit the block limit.

Controlling registers and asking the runtime

Two tools control register use. The -maxrregcount option to nvcc caps registers for every kernel in a file. The __launch_bounds__ qualifier does it per kernel. The Best Practices Guide recommends always giving at least the single argument version, the largest block size the kernel will be launched with, both for forward compatibility and to avoid "too many resources requested for launch" errors. The two argument version adds a minimum number of resident blocks per SM, which lets the compiler choose a register budget that meets it.

// Follows the CUDA C++ Best Practices Guide and Programming Guide.
// Not executed: written on a Mac without an NVIDIA GPU.

// Ask the compiler for a register budget that allows at least
// 4 resident blocks of 256 threads per SM.
// On A100 or H100: 4 blocks x 8 warps = 32 warps, so at most
// 2,048 registers per warp, or 64 registers per thread.
__global__ void __launch_bounds__(256, 4)
scale_kernel(const float* __restrict__ in, float* __restrict__ out,
             float alpha, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = alpha * in[i];
    }
}

The runtime can also do the arithmetic for you at run time. The Best Practices Guide points to cudaOccupancyMaxActiveBlocksPerMultiprocessor, and the runtime also offers cudaOccupancyMaxPotentialBlockSize, which suggests a block size that maximizes occupancy for a given kernel.

// Not executed. Uses the CUDA runtime occupancy API.
#include <cstdio>
#include <cuda_runtime.h>

int main() {
    int device = 0;
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, device);

    int block_size = 256;
    size_t dynamic_smem = 0;
    int blocks_per_sm = 0;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &blocks_per_sm, scale_kernel, block_size, dynamic_smem);

    int active_warps = blocks_per_sm * block_size / prop.warpSize;
    int max_warps = prop.maxThreadsPerMultiProcessor / prop.warpSize;
    printf("SMs: %d, blocks/SM: %d, theoretical occupancy: %.1f%%\n",
           prop.multiProcessorCount, blocks_per_sm,
           100.0 * active_warps / max_warps);

    int min_grid = 0, best_block = 0;
    cudaOccupancyMaxPotentialBlockSize(&min_grid, &best_block,
                                       scale_kernel, dynamic_smem, 0);
    printf("Suggested block size: %d, minimum grid for full GPU: %d\n",
           best_block, min_grid);
    return 0;
}

To see how many registers the compiler actually chose, compile with -Xptxas -v and read the line that reports registers, shared memory and spill stores for each kernel. If you see spill stores and loads, the cap is tighter than the kernel's live state, and you have bought occupancy with local memory traffic.

When occupancy does not matter

A high occupancy figure feels like progress, but it is only one way to give the scheduler independent work. The other is instruction level parallelism inside each warp. Both feed the same Little's law budget. Vasily Volkov's 2016 doctoral thesis at UC Berkeley, "Understanding Latency Hiding on GPUs", studies this trade in detail, and the Best Practices Guide says it plainly: "Higher occupancy does not always equate to higher performance-there is a point above which additional occupancy does not improve performance."

ILP: more work per warp

Compare two versions of a reduction style loop. In the first, each iteration's addition depends on the previous one, so a single warp can have only one addition in flight at a time.

// Not executed. One dependent chain per thread.
__global__ void sum_one_chain(const float* x, float* out, int n) {
    float acc = 0.0f;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;
         i += gridDim.x * blockDim.x) {
        acc += x[i];                  // each add waits for the last
    }
    atomicAdd(out, acc);
}

In the second, each thread keeps four independent accumulators and issues four loads per iteration. The scheduler can issue the second load while the first is still in flight, so the same warp contributes four times as many bytes in flight.

// Not executed. Four independent chains per thread.
__global__ void sum_four_chains(const float4* x4, float* out, int n4) {
    float a0 = 0.f, a1 = 0.f, a2 = 0.f, a3 = 0.f;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n4;
         i += gridDim.x * blockDim.x) {
        float4 v = x4[i];             // one 16-byte load per thread
        a0 += v.x; a1 += v.y; a2 += v.z; a3 += v.w;
    }
    atomicAdd(out, (a0 + a1) + (a2 + a3));
}

The second version needs a few more registers per thread, which can lower occupancy, and it can still be faster, because each warp now moves 512 bytes per load instruction. In our Little's law table, that cut the warps needed per H100 SM from about 54 to about 14. The point is not that this specific kernel will win on your GPU; we have not measured it. The point is that occupancy and ILP trade against each other through the register file, and the right balance is an empirical question.

Register tiling and tensor core kernels

Dense matrix multiply is the extreme case. A well tuned GEMM keeps a tile of the output matrix in registers (or, on Blackwell data center parts, in Tensor Memory), reuses every loaded value many times, and keeps the tensor cores fed from shared memory through asynchronous copies. Such kernels commonly use close to the 255 register maximum per thread. By our table, 128-thread blocks at 255 registers reach only 12.5 percent occupancy, 8 warps per SM, 2 per partition. That is enough when each of those warps has many independent tensor core and load instructions in flight, which is exactly what wgmma.mma_async and TMA were designed to allow on Hopper.

The arithmetic intensity explains why. Multiplying two N by N matrices in FP16 does 2 times N cubed floating point operations and, at minimum, reads and writes 3 times N squared values of 2 bytes. For N equal to 4,096, that is about 1,365 FLOPs per byte, far above what is needed to make the kernel compute bound, as the next section shows.

When low occupancy really hurts

Low occupancy does hurt when each warp has little independent work and the kernel is memory bound: streaming copies, element-wise operations, gathers and scatters. It also hurts when warps spend much of their time at __syncthreads(), because a block with all warps parked at a barrier contributes nothing to the eligible pool. For those kernels, raising occupancy is often the first and cheapest fix.

Rooflines: where the SM meets memory

The roofline model asks whether a kernel is limited by compute or by memory. Divide a GPU's peak FLOP rate by its peak memory bandwidth to get the ridge point, in FLOPs per byte. A kernel whose arithmetic intensity is below the ridge is memory bound; above it, compute bound. All figures below are the vendor's peak numbers from the whitepaper tables, not sustained ones.

A single precision SAXPY, y = a * x + y, does 2 FLOPs per element and moves 12 bytes (two reads and one write of 4 bytes), for an intensity of about 0.17 FLOP per byte. On H100 that caps it at roughly 559 GFLOPS even at full bandwidth, under 1 percent of the FP32 peak. No amount of occupancy changes that ceiling; only less data movement does. The 4,096 square GEMM at 1,365 FLOPs per byte sits well above every ridge point, which is why it can run near peak at low occupancy.

Nsight Compute draws these rooflines for you in its Speed Of Light sections, including hierarchical versions for single, double, half precision and tensor core work.

Reading Nsight Compute metrics

Nsight Compute is NVIDIA's kernel profiler. The sections that answer occupancy and scheduling questions are Launch Statistics, Occupancy, Scheduler Statistics and Warp State Statistics, alongside the Speed Of Light summary. A typical command line, following the Nsight Compute documentation, collects them for one kernel:

# Not executed here. Collects the sections discussed below for one kernel.
ncu --section SpeedOfLight \
    --section LaunchStats \
    --section Occupancy \
    --section SchedulerStats \
    --section WarpStateStats \
    --kernel-name scale_kernel \
    --launch-count 1 \
    -o scale_report ./my_app

The four questions, in order

1. What did the launch ask for? Launch Statistics shows block size, grid size, registers per thread, shared memory per block and the number of waves. A surprise register count shows up here first.

2. What could it achieve? The Occupancy section gives theoretical occupancy and the limit from each resource, backed by metrics you can query directly: launch__occupancy_limit_registers, launch__occupancy_limit_shared_mem, launch__occupancy_limit_warps and launch__occupancy_limit_blocks. The smallest is the binding limit. Metrics such as launch__occupancy_per_register_count and launch__occupancy_per_block_size show how active warps would change if you changed one input.

3. What did it achieve? The guide notes that "large discrepancies between the theoretical and the achieved occupancy during execution typically indicates highly imbalanced workloads." A grid too small to fill all SMs, or a long last wave, also shows up as a gap.

4. Were the schedulers busy? Scheduler Statistics reports theoretical, active and eligible warps per scheduler and how often one issued. In the guide's words, "Having many skipped issue slots indicates poor latency hiding." If eligible warps per scheduler is well above one and few slots are skipped, occupancy is not your problem, however low it looks.

Stall reasons

Warp State Statistics and the source view break the time warps spend not issuing into stall reasons. The ones you see most often, with the Nsight Compute guide's definitions paraphrased:

The full metric names carry a prefix, for example smsp__pcsamp_warps_issue_stalled_long_scoreboard for the sampled version. The smsp prefix tells you the metric is counted per sub-partition, the partition we have been describing all along.

A reading of a hypothetical report

Here is how we would read a report with these numbers, which are invented for illustration: theoretical occupancy 50 percent, limited by registers; achieved 47 percent; 0.4 eligible warps per scheduler; most issue slots skipped; top stall long scoreboard.

The workload is balanced, since achieved is close to theoretical. Fewer than one eligible warp per scheduler, with the dominant stall on global memory, means memory latency is not covered. Either lower registers with __launch_bounds__ to reach 75 or 100 percent occupancy and check for spills, or keep the registers and add independent loads per thread, such as float4 loads. Profile both. Had the report shown 3 eligible warps per scheduler and math pipe throttle on top, we would leave occupancy alone.

A tuning workflow, step by step

1. Baseline. Profile once with the sections above and note duration, achieved occupancy, eligible warps per scheduler and the top two stall reasons.
2. Roofline. If the Speed Of Light section shows memory or compute throughput near peak, scheduling changes will not buy much.
3. Resources. Compile with -Xptxas -v and check registers, shared memory and spills. Register steps fall every 8 registers per thread, because 256 registers per warp divided by 32 threads is 8.
4. One change at a time. Block size, a register cap, a tile size or ILP per thread, then profile again against the baseline.
5. Stop at saturation. When the pipe that bounds the kernel is near peak, further occupancy work will not pay.

The SM and the software that runs on it

Everything in this post runs through a stack of NVIDIA software: the kernel mode driver, the CUDA runtime and toolkit, the PTX compiler inside the driver, and libraries such as cuBLAS, cuDNN and NCCL. In production these usually arrive inside container images, model serving frameworks and Python wheels that bundle their own copies of CUDA libraries. Each of those components has versions, and some versions have published vulnerabilities and security bulletins that require an update.

Safeguard does not design, program or speed up GPUs, and nothing in this post depends on it. Where it fits is the software supply chain around the GPU. Safeguard builds SBOMs and scans repositories, container images and SBOM files (safeguard scan --dir ., safeguard scan --image <name>, safeguard scan --sbom <file>), matches the components it finds against vulnerability databases, and shows which of your projects contain a given component and version. That helps answer questions such as which of our serving images still carry an old library version after a vendor bulletin. It does not cover a GPU driver installed directly on a host outside any image you scan, and the fix for a driver or firmware issue still has to come from NVIDIA or your platform vendor. For more on how accelerators fit together, see our posts on AI chip architecture, what an AI accelerator is and how it works, how a GPU is made, from wafer to HBM package, and estimating GPU cost per million tokens.

Common mistakes

• Treating 100 percent occupancy as the goal. If eligible warps per scheduler is above one and the target pipe is busy, more warps change nothing.
• Dividing 65,536 by the thread count and stopping. That ignores the 256 register per warp rounding and the per-partition banks.
• Forgetting the 1 KB reservation. A block that asks for 48 KB of shared memory consumes 49 KB.
• Treating 48 KB as the hardware limit. It is the static declaration limit. Larger tiles need dynamic shared memory and the cudaFuncSetAttribute opt-in.
• Choosing 1,024-thread blocks on 48-warp parts. On compute capability 8.6, 8.9 and 12.x, one 32-warp block leaves 16 warp slots that no second block can use.
• Capping registers without reading the spill report. A tighter cap can raise occupancy and add local memory traffic that costs more than it saves.
• Using tiny blocks. 32 or 64-thread blocks hit the resident block limit before the warp limit.
• Comparing peaks as if they were sustained. Whitepaper TFLOPS and bandwidth figures are peaks at boost clock.
• Porting tuning between generations blindly. The same kernel can be shared memory bound on A100 and register bound on H100.

FAQ

What is a streaming multiprocessor in an NVIDIA GPU?

It is the GPU's basic processor. Each SM runs many thread blocks at once, holds their registers and shared memory, and schedules their warps. Since Volta it has four partitions, each with its own warp scheduler, register file slice and math units, sharing one L1 and shared memory array. A100 enables 108 SMs and H100 SXM5 enables 132.

How many warp schedulers does an SM have?

Four, one per partition, on Volta, Turing, Ampere and Hopper according to the whitepapers. Each can issue for one warp per clock. The Nsight Compute guide also refers to the four warp schedulers on fifth generation tensor core parts.

How do I calculate occupancy?

Take the minimum of four limits: warps per SM divided by warps per block, the block limit, the blocks that fit in the register file after rounding each warp up to a multiple of 256 registers, and the blocks that fit in shared memory after adding 1 KB per block. Divide the resulting warps by the maximum warps per SM. Or call cudaOccupancyMaxActiveBlocksPerMultiprocessor, or read the Occupancy section in Nsight Compute.

Is higher occupancy always better?

No. NVIDIA's own guides say higher occupancy does not always mean higher performance, though low occupancy always reduces the ability to hide latency. Kernels with many independent instructions per thread, such as tiled matrix multiplies, often run best at low occupancy.

What is the difference between theoretical and achieved occupancy?

Theoretical occupancy is what the launch configuration allows on paper. Achieved occupancy is the measured average of active warps during the run. A large gap usually means an imbalanced workload, a grid too small to fill the GPU, or a long last wave.

How much shared memory can one block use on H100?

Up to 227 KB, out of 228 KB of shared memory per SM, because CUDA reserves 1 KB per block. Anything above 48 KB must be dynamic shared memory with an explicit opt-in. On A100 the figures are 163 KB per block and 164 KB per SM.

What is the maximum number of registers per thread?

255, on every compute capability from 7.5 to 12.x. At 128 threads per block, 255 registers limits A100 and H100 to 12.5 percent occupancy by our calculation, which tiled matrix kernels often accept.

What does the long scoreboard stall mean in Nsight Compute?

The warp is waiting for data from an L1TEX operation, meaning a local, global, surface or texture memory access. It is the classic sign that memory latency is not covered. Increase bytes in flight per warp, improve coalescing, or add warps.

The SM itself is hardware, but the code that drives it is software you install and update: drivers, the CUDA toolkit, cuDNN, NCCL and the serving stacks and container images that bundle them. Those components carry open source dependencies and get security bulletins like any other software, so the same patching discipline applies to the images you build around your GPUs.

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
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• 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
• 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 the CLI, open the interactive terminal and ask it to fix a lodash CVE in a repository; it reads the finding and the advisory, asks before any command that changes something, edits package.json, runs the install and tests with your approval and pushes nothing. The same CLI's /scan . and safeguard scan --image commands are how you would inventory a model serving repository or the container image it ships in.

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

The second demo sets up the Safeguard runner on your own machine, approves it with a device code and then scans from the terminal. That is the shape you would use on a GPU build host, where images are built and scanned on hardware you control before they reach a cluster.

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 Safeguard tools Claude can use. Once connected, you can ask in plain language which projects contain a given library version, for example after a GPU library security bulletin.

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

To try it on your own repositories and images, download the Safeguard CLI and follow the CLI documentation.

Quick check

Sources

• NVIDIA: CUDA Programming Guide v13.4.2, Compute Capabilities appendix (feature and technical specification tables)
• NVIDIA: CUDA Programming Guide, Writing CUDA Kernels (occupancy, maxrregcount)
• NVIDIA: CUDA Programming Guide, Advanced Kernel Programming (independent thread scheduling, async proxy, carveout)
• NVIDIA: CUDA Programming Guide, Programming Model
• NVIDIA: CUDA Programming Guide, Understanding Memory
• NVIDIA: CUDA C++ Best Practices Guide (register allocation rounding, occupancy example, latency)
• NVIDIA: Volta Tuning Guide
• NVIDIA: Ampere Tuning Guide
• NVIDIA: Hopper Tuning Guide
• NVIDIA: Blackwell Tuning Guide
• NVIDIA: Parallel Thread Execution ISA v9.4 (warpgroups, wgmma, Tensor Memory)
• NVIDIA: Nsight Compute Profiling Guide v2026.3.1 (occupancy, scheduler statistics, warp states, stall reasons, metrics)
• NVIDIA: Tesla V100 GPU Architecture whitepaper (Volta)
• NVIDIA: Turing GPU Architecture whitepaper
• NVIDIA: A100 Tensor Core GPU Architecture whitepaper (Ampere GA100)
• NVIDIA: Ampere GA102 GPU Architecture whitepaper v2
• NVIDIA: H100 Tensor Core GPU Architecture whitepaper V1.04 (Hopper)
• Vasily Volkov: Understanding Latency Hiding on GPUs, UC Berkeley EECS technical report EECS-2016-143 (PhD thesis, 2016)

Last reviewed: 25 October 2026.