In 2010 Vasily Volkov’s GTC talk “Better Performance at Lower Occupancy” argued that the GPU community is wrong to focus on occupancy as it doesn’t directly relate to throughput. Throughput often rises as occupancy falls because what really keeps a GPU busy is per thread register budget and instruction level parallelism, not how many warps you cram into an SM. This blog re-examines each of his claims on modern hardware (fp32 cuda cores + registers) and reports what survived and what didn’t.

Occupancy


I use RTX 2000 Ada for the experiments done here. Also, since this run is on a RunPod container where GPU performance counters are admin locked, instead of NCU’s achieved occupancy I use theoretical occupancy (from cudaOccupancyMaxActiveBlocksPerMultiprocessor).


Occupancy vs. Throughput

I wrote a compute bound fp32 GEMM to test Volkov’s occupancy claim on modern hardware. The idea is to hold the math per output block constant and vary only the sub tile per thread. A bigger sub tile means fewer threads share the block, more live accumulators per thread, more registers and lower occupancy. I fixed the output block at 64×64 and only let the subtile size vary. I used __launch_bounds__(threads) deliberately so the register count stays determinant and is driven by actual registers required, not compiler’s guess (see here). I timed each variant for 50 iterations using cudaEvent timing and plotted the theoretical occupancy (on x-axis) vs achieved GFLOP/s (on y-axis). Check the code in the Appendix.

Occupancy vs Throughput

Variant Threads/block Sub-tile TM×TN Regs/thread Theo. occ GFLOP/s
S 1024 2×2 62 66.7% 1394
A1 512 2×4 40 100.0% 2651
A 256 4×4 64 66.7% 3770
B 128 8×4 96 41.7% 4230
C 64 8×8 150 25.0% 3713
D 32 16×8 254 16.7% 2970
D+ 16 16×16 255 8.3% 411

The conventional model of “more concurrent warps = more latency hiding = more throughput” predicts a monotonic curve with 100% occupancy variant (A1) on the top and the 8% occupancy variant (D+) at the bottom. In practice the peak is at 42% while D+ collapses for another reason (register spill).

The cause is per thread ILP not occupancy. Large sub tile packs more live accumulators into registers giving each warp independent FMAs to interleave without warp switching. The proof is S and A sitting at the same 66.7% occupancy while A is 2.7× faster. Occupancy is the symptom and register driven ILP is the cause. Past B the curve falls for a reason Volkov didn’t discuss, D at 254 registers, leaves the compiler no scheduling headroom and D+ pushes past the 255 register cap collapsing to 411 GFLOP/s (10× lower than the peak) as the compiler spills to local memory (covered in section 3). No variant nears the 12 TFLOP/s ceiling but that’s expected as this is a hand written kernel with no tensor cores or double buffering.


Memory bandwidth at low occupancy

For testing “you need lots of threads to hide memory latency”, I wrote a memcpy kernel. By sweeping occupancy from 67% down to about 4% and watching whether achieved memory bandwidth collapses or stays saturated.

To isolate what drives bandwidth I use two knobs, thread count and bytes-in-flight. Each variant is a strided copy with __launch_bounds__ and a 48 KB per block dynamic shared-memory allocation caps blocks/SM at 2. Volkov suggested to allocate smem to control residency. What I sweep is loads × fetch width per thread. loads is how many independent load instructions each thread issues (its latency hiding depth) and width is the fetch width of each (float=1, float2=2 and float4=4). So the sweep drops thread count by 16× while raising bytes-in-flight per thread 56× (4 to 224). This is not a tile, section 1’s TM × TN was about per thread compute/accumulators while loads × width is about memory transaction shaping.

I ran each variant on a 256 MB buffer for 50 iterations. Peak memory on this card is 224 GB/s and nvidia-smi dmon (sampled at 1 Hz during the run) pinned the mem utilization column at 100% throughout. Check the code in the Appendix.

Bandwidth vs Occupancy

Variant Threads/block Loads × width Bytes/thread Regs/thread Theo. occ GB/s
T1 1024 1×1 4 8 66.7% 194.7
T2 1024 2×1 8 11 66.7% 197.9
T4 512 4×1 16 21 66.7% 206.7
T8 256 8×1 32 38 33.3% 206.7
T8f2 128 8×2 64 38 16.7% 206.7
T8f4 64 8×4 128 40 8.3% 205.7
T14f4 32 14×4 224 50 4.2% 199.3

Note that Regs/thread is the compiler’s total register cost which includes the two 64 bit base pointers (src and dst) at 2 registers each plus loop index, stride and bound. The payload per thread (loads × width × 4 bytes) is what keeps memory busy.

The folklore says GPU memory is slow (hundreds of cycles) so you need lots of threads to hide the latency. Maximize occupancy or the bandwidth collapses. The experiment contradicts this. Bandwidth holds at about 92% of memory peak all the way from 67% down to 4% occupancy. A 16× drop in thread count costs just 3% of bandwidth. What actually keeps memory busy is bytes in flight not threads (Little’s Law). One thread issuing 14 float4 loads keeps as much in flight as 56 threads each issuing one float. So occupancy (a thread count) is the wrong dial, bytes-in-flight is the right one. Give each thread enough outstanding loads and you don’t need many threads.


Register pressure: the modern ceiling

Section 1 lowered occupancy by giving each thread more useful work (bigger tile leading to more accumulators leading to more registers). Here I do the opposite I take one kernel (Section 1’s variant C) and artificially cap its register count with -maxrregcount=N. Forcing the compiler to spill whenever N is below what the kernel wants and ask where exactly the spill cliff sits.

I recompiled variant C (TM=TN=8) eight times with -maxrregcount ∈ {32, 48, 64, 96, 128, 160, 192, 255}. Each binary’s -Xptxas -v log gives me the actual register count, spill stores, spill loads, and stack frame. I then time each at 4096³, 50 iters with only C. No profiler is used here. The spill numbers come straight from the compiler. ( You can check the code in the Appendix)

Register Pressure vs Occupancy

Cap (-maxrregcount) Actual regs Theo. occ ms/call GFLOP/s Spill stores Spill loads Stack frame
32 32 41.7% 547.44 251 4840 4652 1672
48 48 41.7% 331.78 414 3160 2964 904
64 64 41.7% 223.81 614 2140 1996 696
96 96 41.7% 41.15 3340 200 140 80
128 128 33.3% 32.47 4233 0 0 0
160 160 25.0% 36.79 3736 0 0 0
192 154 25.0% 36.92 3723 0 0 0
255 154 25.0% 36.86 3729 0 0 0

This is a cliff, not a curve. Below 128 regs the kernel spills thousands of bytes to local memory and throughput crawls from 251 up to 3340 GFLOP/s. The 32 reg variant is 17× slower than peak. At 128 regs the smallest spill free budget spills hit zero and throughput jumps to 4233 GFLOP/s. Above the cliff more registers hurt. When going from 128 to 255, throughput drifts back down to 3730 because the compiler stops allocating at 154 actual regs. So the extra budget just lowers occupancy (going from 8 to 6 blocks/SM) without buying ILP. The “natural” register count isn’t always the best one. So Volkov’s “use more registers” still holds with one caveat that we should spend registers to buy ILP but never past the spill cliff. The same variant C that in section 1 gave 3713 GFLOPS/s hits 4233 at 128 capped registers.

What survived

The thread across all three sections is the same, occupancy is the wrong dial. For compute the real lever is per thread ILP through the register file (section 1), bounded above by the spill cliff (section 3). For memory it’s bytes-in-flight, not thread count (section 2). Volkov was right in 2010 and he’s still right on modern fp32 hardware with one footnote, the register lever is a hard ceiling and overshooting it is worse than not using it at all.

All three sections ran on fp32 CUDA cores but a modern GPU does most of its throughput through tensor cores and that regime shift changes both levers. In Part II I ask what tensor cores do to the register lever and whether the modern warp scheduler has absorbed the ILP lever too.

Appendix

Below are the kernel-only code. The host side harness (launch setup, timing, the template param / -maxrregcount sweep) is elided for brevity.

Occupancy vs Througput code

template <int BM, int BN, int BK, int TM, int TN>                   // template for different shapes
__global__ __launch_bounds__((BM * BN) / (TM * TN))                 // (64*64)/(4*4) - threads per block
void gemm_kernel(const float* __restrict__ A,
                 const float* __restrict__ B,
                 float* __restrict__ C,
                 int M, int N, int K, float alpha, float beta) {
    constexpr int THREADS = (BM * BN) / (TM * TN);
    constexpr int COLS = BN / TN;

    const int bx = blockIdx.x;  // along N
    const int by = blockIdx.y;  // along M
    const int tx = threadIdx.x;
    const int tr = tx / COLS;
    const int tc = tx % COLS;

    __shared__ float As[BM][BK + 1];   // +1 pad to reduce bank conflicts on BK stride
    __shared__ float Bs[BK][BN + 1];
    
    // set the accumulator
    float acc[TM][TN];
    #pragma unroll
    for (int m = 0; m < TM; m++)
        #pragma unroll
        for (int n = 0; n < TN; n++) acc[m][n] = 0.0f;
    
    // tiled matrix mul
    #pragma unroll 1
    for (int kk = 0; kk < K; kk += BK) {
        #pragma unroll
        for (int i = tx; i < BM * BK; i += THREADS) {
            const int r = i / BK;
            const int c = i % BK;
            As[r][c] = A[(by * BM + r) * K + (kk + c)];
        }
        #pragma unroll
        for (int i = tx; i < BK * BN; i += THREADS) {
            const int r = i / BN;
            const int c = i % BN;
            Bs[r][c] = B[(kk + r) * N + (bx * BN + c)];
        }
        __syncthreads();

        #pragma unroll
        for (int k = 0; k < BK; k++) {
            float ar[TM];
            #pragma unroll
            for (int m = 0; m < TM; m++) ar[m] = As[tr * TM + m][k];
            float br[TN];
            #pragma unroll
            for (int n = 0; n < TN; n++) br[n] = Bs[k][tc * TN + n];

            #pragma unroll
            for (int m = 0; m < TM; m++)
                #pragma unroll
                for (int n = 0; n < TN; n++)
                    acc[m][n] = ar[m] * br[n] + acc[m][n];
        }
        __syncthreads();
    }

    #pragma unroll
    for (int m = 0; m < TM; m++)
        #pragma unroll
        for (int n = 0; n < TN; n++) {
            const int row = by * BM + tr * TM + m;
            const int col = bx * BN + tc * TN + n;
            if (row < M && col < N)
                C[row * N + col] = alpha * acc[m][n] + beta * C[row * N + col];
        }
}

Bandwidth vs Occupancy

template <int THREADS, int V, int WT>
__global__ void __launch_bounds__(THREADS)
memcpy_kernel(const float* __restrict__ src, float* __restrict__ dst, size_t n) {
    extern __shared__ char occ_smem[];
    if (threadIdx.x == 0) occ_smem[0] = (char)0;

    // strided grid: thread t (global index g) writes contiguous floats,
    // but iterates over V tiles separated by gridStride.  This issues V
    // independent loads per thread so the memory subsystem can keep them
    // all in flight (Volkov's "no latency stall until dependent op").
    const size_t grid_stride = (size_t)gridDim.x * THREADS;
    size_t tid = (size_t)blockIdx.x * THREADS + threadIdx.x;
    for (int i = 0; i < V; ++i) {
        size_t base = tid + (size_t)i * grid_stride;
        if (WT == 4) {
            const float4* s4 = reinterpret_cast<const float4*>(src);
            float4*       d4 = reinterpret_cast<float4*>(dst);
            size_t idx = base;  // floats already in float4 units when WT==4
            if (idx * 4 + 3 < n)
                d4[idx] = s4[idx];
        } else if (WT == 2) {
            const float2* s2 = reinterpret_cast<const float2*>(src);
            float2*       d2 = reinterpret_cast<float2*>(dst);
            size_t idx = base;
            if (idx * 2 + 1 < n)
                d2[idx] = s2[idx];
        } else {
            if (base < n) dst[base] = src[base];
        }
    }
}

Spill vs Occupancy

Same as the first kernel in Occupancy vs Throughput