100 DAYS OF GPU Course preview
Day 45Module 5
in-technical-review

Occupancy is not the goal

Two earlier measurements challenge the claim that more occupancy is always better.

On day 17 a kernel ran at 75 percent occupancy in 0.164 ms. Adding __launch_bounds__ took it to 100 percent and it slowed to 0.227. On day 10 cudaOccupancyMaxPotentialBlockSize suggested 1024 threads a block, 1024 was not the fastest row, and nine of thirteen block sizes landed within 5 percent of the best one.

Both results make sense because occupancy does not measure speed. This page tunes one kernel two ways to reach the same bandwidth at occupancies that differ by a factor of four. Four profiler rows show which resource limits each configuration.

What the warp slots are for

An SM holds a fixed number of warps at once. Day 2 showed how to read the warp and block limits for a GPU.

Occupancy is how many of those warp slots your kernel fills. Nsight Compute puts it in one line: "Occupancy is the ratio of the number of active warps per multiprocessor to the maximum number of possible active warps" (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html , checked 2026-08-30).

Warp slots support latency hiding. A warp that issues a load cannot use the result for hundreds of cycles, so the scheduler sets it aside and issues from a different warp. With enough warps resident the memory system never goes quiet.

With too few, memory requests pause and the kernel waits on DRAM.

The same page is careful about which way the implication runs: "Higher occupancy does not always result in higher performance, however, low occupancy always reduces the ability to hide latencies, resulting in overall performance degradation" (same page, checked 2026-08-30). Low occupancy can limit latency hiding, but higher occupancy does not by itself mean higher speed.

Four resources limit how many blocks fit on an SM. Warp slots limit blocks because a block of T threads needs ceil(T / 32) warps. The device also has a fixed number of block slots per SM.

The register file sets another limit because each block needs its threads times registers per thread. Shared memory works the same way: asking for more per block leaves room for fewer blocks.

For the measured T4, a 256-thread block uses 8 of 32 warp slots. Each SM has 16 block slots, 65,536 registers, and 64 KiB of shared memory. These values describe that result, not general CUDA limits.

Nsight Compute reports all four as separate rows, named Block Limit Warps, Block Limit SM, Block Limit Registers and Block Limit Shared Mem, with the metrics launch__occupancy_limit_warps, launch__occupancy_limit_blocks, launch__occupancy_limit_registers and launch__occupancy_limit_shared_mem (same page, checked 2026-08-30). Those four rows answer "why is my occupancy not 100 percent". They identify the resource that limits each launch.

The register-bound preset starts with the register file as the limit. Drag the shared memory control to change which resource limits residency. The widget models one SM and does not predict run time.

Occupancy does not count requests

Occupancy counts warps. It does not count the requests each warp has in flight.

Little's law connects concurrency, latency, and throughput. Concurrency is latency times throughput, so a busy memory system needs a certain number of requests in flight. Warps are only one way to supply them, and one warp with eight independent loads outstanding does the same job as eight warps with one each.

That is instruction-level parallelism, and it uses registers instead of warp slots.

Resident warps and independent requests per warp both add concurrency. More requests per warp need more registers, which can reduce the number of resident warps. Register pressure names this limit, and day 17 measured its cost.

More threads are not the only source of parallel work on a GPU. A thread may have one request in flight or eight. The loads must issue before the thread uses the first result:

// load, use, store per item: one request in flight, whatever `items` is,
// because the multiply on line three waits for the load on line three
for (int j = 0; j < items; ++j) {
    const size_t i = base + j * step;
    out[i] = a * x[i] + y[i];
}

Ten configurations, three settings

Full program in code/day45-occupancy/occupancy.cu. Ten rows, one kernel body, and every row moves the same 192 MiB.

The block size never moves. 256 threads on every row. Day 10 swept that and held everything else still; this page does the reverse, because a table where two things changed cannot attribute a result to either.

One setting changes occupancy and nothing else. The rows labelled -bN reserve dynamic shared memory that the kernel never declares and never reads, so that only N blocks fit on an SM. The instructions are identical to the row above; only the third launch argument differs. This isolates occupancy from the kernel's work.

Residency comes from the driver, not from arithmetic on this page. cudaOccupancyMaxActiveBlocksPerMultiprocessor answers about the compiled kernel, so it accounts for a register count the source never shows you, and it takes the shared-memory reservation as an argument.

        CUDA_CHECK(cudaFuncGetAttributes(&attrs, scaleAdd<kItems>));
        CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
            &blocksPerSm, scaleAdd<kItems>, kThreadsPerBlock, smem));

The program checks correctness at every row, because the table depends on the claim that the kernel covers the array for any grid and any items-per-thread setting. One check at one row would not test that claim. The program zeros the output buffer first and returns EXIT_FAILURE on an error instead of using an assert, which NDEBUG would remove.

The two loops in the body are separate on purpose, and that separation is the experiment: every load issues before the first multiply, so one warp holds 2 * kItems requests at once.

template <int kItems>
__device__ void scaleAddBody(float a, const float* __restrict__ x,
                             const float* __restrict__ y,
                             float* __restrict__ out, size_t n) {
    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    const size_t base =
        blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;

    float xv[kItems];
    float yv[kItems];
#pragma unroll
    for (int j = 0; j < kItems; ++j) {
        const size_t i = base + static_cast<size_t>(j) * step;
        xv[j] = (i < n) ? x[i] : 0.0f;
        yv[j] = (i < n) ? y[i] : 0.0f;
    }
#pragma unroll
    for (int j = 0; j < kItems; ++j) {
        const size_t i = base + static_cast<size_t>(j) * step;
        if (i < n) {
            out[i] = a * xv[j] + yv[j];
        }
    }
}

#pragma unroll with a literal j is what keeps xv and yv in registers. Give the index a runtime value and the arrays become local memory, which day 17 measured costing 9.5x.

Host code changes the occupancy setting without changing the kernel:

static size_t sharedForBlocks(const cudaDeviceProp& prop, int target) {
    if (target < 1) {
        return 0;
    }
    size_t want = prop.sharedMemPerMultiprocessor / static_cast<size_t>(target);
    want &= ~static_cast<size_t>(255);
    if (want > prop.sharedMemPerBlock) {
        want = prop.sharedMemPerBlock;
    }
    return want;
}

The third setting is a budget, not a description of the launch: "the compiler first derives the upper limit, L, on the number of registers that the kernel should use ... If the initial register usage exceeds L, the compiler reduces it until it is less than or equal to L. This usually results in increased local memory usage and/or a higher number of instructions" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html section 5.4.3.2, checked 2026-08-30).

template <int kItems>
__global__ __launch_bounds__(kThreadsPerBlock, kMinBlocksPerSm) void
scaleAddBounded(float a, const float* __restrict__ x,
                const float* __restrict__ y, float* __restrict__ out,
                size_t n) {
    scaleAddBody<kItems>(a, x, y, out, n);
}

Note. The shared-memory reservation is a residency knob and nothing else. The kernel declares no extern __shared__ array and reads none, so a row and its -bN sibling run the same instructions on the same addresses. The one thing the reservation does not model is a kernel that pays for shared memory by using it, which is day 13.

Results

Measured on a Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88), built with nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo, timed with CUDA events, mean of 10 runs after 3 warm-ups, 192.0 MiB moved per pass in every row.

Captured 2026-09-02 on the project's verification node; the transcript and fresh Nsight Systems capture are in evidence. The ten per-configuration Nsight Compute reports at content/profiles/occupancy/ remain the CUDA 12.6, driver 595.84 counter evidence captured on 2026-09-01.

label items a thread reservation occupancy in flight ms GB/s
tlp1 one none 100% 64 0.770 261.4
tlp1-b3 one three blocks an SM 75% 48 0.823 244.7
tlp1-b2 one two blocks an SM 50% 32 1.107 181.9
tlp1-b1 one one block an SM 25% 16 2.001 100.6
ilp2 two none 100% 128 0.775 259.9
ilp4 four none 100% 256 0.801 251.3
ilp8 eight none 100% 512 0.810 248.4
ilp4-b1 four one block an SM 25% 64 0.792 254.2
ilp8-b1 eight one block an SM 25% 128 0.791 254.4
ilp8-lb eight none, plus a launch bound 100% 512 0.810 248.6

The three claims and the measurements

Held. tlp1 took 0.770 ms and ilp4-b1 took 0.792, a ratio of 1.03, at occupancies four times apart. Sixty-four loads in flight per SM is sixty-four loads in flight, whether thirty-two warps carry two each or eight warps carry eight. Little's law was the right model, and the program prints that pair as its last line.

Held, with a larger low-ILP drop. tlp1-b1 is the slowest row at 2.001 ms, 0.38 of the best, and the occupancy gap alone cannot be why: ilp4-b1 and ilp8-b1 sit at the same 25 percent occupancy and run within 3 percent of the fastest row. Under CUDA 12.6, tlp1-b2 and tlp1-b1 reached 254.3 and 176.7 GB/s; under CUDA 13 they fell to 181.9 and 100.6 GB/s while the ILP-protected 25-percent rows held above 254 GB/s. The drift strengthens the mechanism: what tlp1-b1 lacks is not occupancy alone, it is the sixteen loads in flight they collectively carry.

Held, with less change than predicted. ilp8-lb and ilp8 are effectively identical, both 0.810 ms and about 248.5 GB/s. The prediction expected day 17's spill mechanism, but local B stayed zero: at 54 registers a thread this kernel already fits four blocks an SM, so __launch_bounds__(256, 4) promised the compiler something it was doing anyway. A launch bound that does not change the register limit has no cost.

The claims did not cover the pure ILP column. It is mildly but consistently slower at full occupancy (0.770, 0.775, 0.801, 0.810 down the tlp1, ilp2, ilp4, ilp8 line). With all thirty-two warp slots full, deeper per-thread pipelining adds no useful requests and needs more address arithmetic.

Per-thread parallelism helps when low occupancy limits the number of requests in flight. It does not help when all warp slots are full.

Compare the order of the rows. The test asks whether occupancy or loads in flight better predicts that order. The exact figures describe the measured GPU.

Run it yourself

The timing half runs on a supported CUDA GPU:

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o occupancy occupancy.cu

Not Compiler Explorer. Three 64 MiB buffers and 140 kernel launches are well past its 20 second run cap, and shrinking the arrays until they fit would shrink the very thing being kept busy.

The profiler half needs counter permission. On this project's node, plain ncu returns ERR_NVGPUCTRPERM, whose full text is The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device. sudo ncu works. The cause on this node is RmProfilingAdminOnly: 1 in /proc/driver/nvidia/params, measured again on 2026-09-02.

The current free-Colab case remains untested. You can finish the exercise from the shipped reports when counters are unavailable, at content/profiles/occupancy/: occupancy-<label>.ncu-rep, one per row of the table, plus occupancy.nsys-rep and a details.csv beside each. Every profiler number this page quotes is in those files as text.

Exercise

Find the lowest-occupancy configuration that still reaches within a tenth of the fastest row, then say in two sentences what is holding its bandwidth up.

Time: 30 to 45 minutes. Submit: your two tables, the settings of the row you found, and the two sentences.

Check: the program compares against a CPU reference at every configuration and returns EXIT_FAILURE naming the first index that disagrees, so a pass means the kernel really covered the array at every launch shape. Four more branches fail the run if the driver's residency breaks a per-SM cap.

It then prints the fastest row and the paired comparison: the two occupancies, the two loads-in-flight figures and the ratio of the times. Report that ratio and the occupancy gap beside it, never the milliseconds, which depend on the GPU.

Hint 1

One setting uses warp slots and the other uses registers. You can change the second setting. How far can you increase it before the row stops getting faster, and what appears in the ceilings table after that point?

Hint 2

Read the in flight column against GB/s, then read occ against GB/s, and see which relation stays consistent. Then look at what regs and local B are doing in the rows where in flight is largest.

Solution

The lowest-occupancy row that still reaches the top is the lowest one whose resident warps times loads per thread is still at least the fastest row's. Below that product, time rises roughly in proportion because too few requests reach the memory system. Above it, extra requests wait because DRAM is already saturated.

Increasing the items per thread past that point also has a cost. More items a thread means more live values, so registers rise, then residency falls, then the compiler puts the array in local memory, which raises memory traffic and shows up as a non-zero local B before it shows up in the time.

Occupancy is one factor in the number of requests in flight. Report the ratio between your rows rather than their absolute times, because the ratio is the finding and the absolute times depend on the GPU.

Pitfalls

You raised occupancy and the kernel got slower. A tight __launch_bounds__ setting limits registers, and ptxas may meet it by spilling. Day 17 measured 100 percent occupancy at 0.227 ms against 75 percent at 0.164 on the same kernel. Build with -Xptxas -warn-spills to make ptxas report spills: ptxas warning : Registers are spilled to local memory in function.

Register spilling is day 17.

You took the occupancy API's answer and shipped it. cudaOccupancyMaxPotentialBlockSize maximises occupancy, which is not the same as minimising time. On day 10 it suggested 1024 threads a block on a memory-bound kernel and 1024 was not the fastest launch configuration measured.

Achieved occupancy is lower than theoretical and you assumed a bug. The theoretical figure is a ceiling the driver will allow; the achieved one is what the kernel reached, and it falls whenever the grid is too small to fill the GPU or blocks retire unevenly at the tail. Day 42 reads the two columns side by side.

The launch fails with too many resources requested for launch. Threads per block times registers per thread exceeded the register file, or you launched more threads per block than the kernel's own __launch_bounds__ allows. The bound's first argument is a hard cap on the launch, not a hint.

You cut registers with -maxrregcount to increase occupancy. That option is file-wide, so it hits every kernel in the translation unit including the ones that were fine. __launch_bounds__ is per kernel. Day 17 has the numbers.

You tuned occupancy on a compute-bound kernel. Occupancy helps hide memory latency. If the SM issues math nearly every cycle, extra warps only reduce the register budget.

Day 30 shows how to tell which limit applies. Day 42's speed-of-light section uses profiler data for the same check, and stall reasons name what the warps wait for.

Go deeper

Next

Day 46 drops below the report into PTX and SASS, where the two loops in this kernel become a visible run of loads before the first multiply and you can count them yourself. Day 48 fuses a four-kernel chain, which is the third way to give a resident warp more to have in flight, and day 49 puts every kernel you have written on one plot, where occupancy does not appear as an axis at all.