2.3 Occupancy and resource allocation

Occupancy is the ratio of active warps on an SM to the theoretical maximum, and one of the core metrics when tuning CUDA kernels. This article covers the definition, how to compute it, the three main limiters (registers / shared memory / block size), and why higher occupancy is not always better — the real goal is a balance between latency hiding and resource use.

📑 Contents

• 1. What is occupancy
• 2. SM resource inventory
• 3. Three factors that limit occupancy
• 4. Occupancy calculation walkthrough
• 5. Higher occupancy is not always better
• 6. Tuning strategies and tools
• 7. Case studies
• Summary
• Self-check list
• References

---

1. What is occupancy

Picture a parking lot with 64 spots (the max warps that can reside on an SM). If only 32 cars are parked, utilization is 50%. Occupancy is that “parking lot utilization” for an SM — how much of the hardware capacity you are actually using.

1.1 Formal definition

$$\text{Occupancy} = \frac{\text{active warps per SM}}{\text{max warps supported per SM}}$$

For example on A100 (Compute Capability 8.0):

• Each SM supports up to 64 warps (2048 threads)
• If your kernel is resource-limited to 32 warps per SM
• Then Occupancy = 32/64 = 50%

1.2 Why occupancy matters

The main value of occupancy is latency hiding. The GPU execution model relies on switching warps to cover memory access latency:

Timeline:
Warp 0: [compute] [wait for memory...400cycles...] [compute]
Warp 1:           [compute] [wait for memory...400cycles...] [compute]
Warp 2:                     [compute] [wait for memory...400cycles...]
...

If enough warps are active, the scheduler can always find a ready warp to fill the wait

How many warps do you need to fully hide latency? A rough estimate:

$$\text{warps needed} \geq \frac{\text{memory latency (cycles)}}{\text{cycles per instruction}} = \frac{400}{4} = 100$$

But an SM tops out at 64 warps, so you cannot fully hide latency in practice — which is why you push occupancy as high as you can (though that is not the only goal).

---

2. SM resource inventory

Each SM has a fixed set of hardware resources shared by every block that resides on it. Knowing the caps is the foundation for computing occupancy.

2.1 Main architecture resource comparison

2.2 Resource allocation hierarchy

graph TD
    A["SM total resource pool"] --> B["Block 0"]
    A --> C["Block 1"]
    A --> D["Block N"]
    B --> E["Warp 0-7"]
    C --> F["Warp 8-15"]
    D --> G["Warp ..."]
    E --> H["Thread 0-255"]

📌 Key point: Resources are allocated in block units. A block either fully enters an SM or not at all — there is no half-block residency. So if one block needs just a bit more than half of the remaining SM resources, the leftover space is wasted.

---

3. Three factors that limit occupancy

3.1 Factor one: register usage

Registers per thread are the most common occupancy limiter.

Computation:

$$\text{resident threads per SM} = \lfloor \frac{\text{total SM registers}}{\text{registers per thread}} \rfloor$$

Actual allocation also has a granularity constraint. Registers are handed to a warp in groups of 256 (allocation granularity). So registers allocated per warp = $\lceil \frac{\text{registers per thread} \times 32}{256} \rceil \times 256$.

Example (A100, 65536 registers/SM):

💡 Tip: You can cap register use with __launch_bounds__ or the maxrregcount compile option:

// Approach 1: hint the compiler with __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
Memory-bound (bandwidth bottleneck)  __launch_bounds__(N, large value)  → raise occupancy
Compute-bound (FLOP bottleneck)      __launch_bounds__(N, 1)            → keep registers

__global__ void __launch_bounds__(256, 2)  // max 256 threads/block, target 2 blocks/SM
my_kernel(...) { ... }

// Compiler behavior
SM register total (A100) = 65536 / SM
maxThreadsPerBlock=256, minBlocksPerMultiprocessor=2
↓
Need concurrent residency: 256 threads × 2 blocks = 512 threads
↓
Max registers per thread: 65536 / 512 = 128
↓
Compiler sets the register cap to 128 (often compresses more aggressively)
↓
Overflow variables spill to Local Memory (device memory)

// Approach 2: global compile-time limit
// nvcc -maxrregcount=32 my_kernel.cu

⚠️ Note: Forcing a lower register count can make the compiler spill variables to local memory (register spilling), which can hurt performance. Check with a profiler.

3.2 Factor two: shared memory usage

The total shared memory declared per block (static + dynamic) limits how many blocks fit on an SM.

Computation:

$$\text{resident blocks per SM} = \lfloor \frac{\text{total SM shared memory}}{\text{shared memory per block}} \rfloor$$

Example (A100, 164KB shared memory/SM):

Shared memory allocation is not byte-granular either; it uses units of 128 Bytes (or 256 Bytes, depending on the architecture).

3.3 Factor three: block size

Block size is limited in two ways:

1. Max 1024 threads per block
2. Max blocks per SM (32 on A100)

An easy-to-miss limit: if your blocks are tiny (e.g. 32 threads = 1 warp), even with low per-block resource use you can still hit the “max blocks per SM” cap:

📌 Key point: Prefer block sizes of at least 128 or 256 so you avoid the block-count limit and give the compiler more room to optimize. 256 is usually a solid default.

---

4. Occupancy calculation walkthrough

4.1 Combined example

Given (A100):

• Kernel uses 40 registers per thread
• Block size: 256 threads = 8 warps
• 16 KB shared memory per block

Step 1: register limit

$$\text{registers per warp} = 40 \times 32 = 1280 \rightarrow \text{align to} 1280$$ $$\text{warps that fit on SM} = \lfloor \frac{65536}{1280} \rfloor = 51 \text{ Warps}$$

Step 2: shared memory limit

$$\text{blocks that fit on SM} = \lfloor \frac{164\text{KB}}{16\text{KB}} \rfloor = 10 \text{ Blocks}$$ $$\text{corresponding warps} = 10 \times 8 = 80 \text{ Warps}$$

Step 3: block-count limit

Max 32 blocks per SM; 10 blocks do not hit the cap.

Step 4: thread-count limit

Max 2048 threads per SM = 64 warps.

Take the minimum:

$$\text{actual warps} = \min(51, 80, 64) = 51 \text{ Warps}$$

But blocks are allocated whole (8 warps each), so in practice:

$$\lfloor \frac{51}{8} \rfloor \times 8 = 6 \times 8 = 48 \text{ Warps}$$ $$\text{Occupancy} = \frac{48}{64} = 75\%$$

4.2 Using the CUDA Occupancy API

#include <cuda_runtime.h>

int main() {
    int blockSize = 256;
    int minGridSize, gridSize;

    // Auto-compute a good block size
    cudaOccupancyMaxPotentialBlockSize(
        &minGridSize, &blockSize,
        my_kernel,  // kernel function pointer
        0,          // dynamic shared memory size
        0           // block size upper bound (0 = no limit)
    );

    // Query occupancy for a given config
    int maxActiveBlocks;
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &maxActiveBlocks,
        my_kernel,
        blockSize,
        0  // dynamic shared memory
    );

    int device;
    cudaGetDevice(&device);
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, device);

    float occupancy = (float)(maxActiveBlocks * blockSize) /
                      prop.maxThreadsPerMultiProcessor;
    printf("Occupancy: %.1f%%\n", occupancy * 100);
    return 0;
}

---

5. Higher occupancy is not always better

One of the most important ideas in CUDA tuning: high occupancy ≠ high performance.

5.1 A counterintuitive example

Consider two versions of a matrix-multiply kernel:

Version B has lower occupancy but higher performance. It uses more registers and shared memory to increase data reuse: each load from global memory is reused many times, so total memory traffic drops.

5.2 The cost of chasing high occupancy

Blindly maximizing occupancy can lead to:

1. Register spilling: forcing a lower maxrregcount makes the compiler spill variables to local memory (really global memory + L1), which can be tens of times slower

2. Weak shared-memory reuse: less shared memory means smaller tiles, lower reuse, and often more global memory traffic

3. Cache thrashing: too many active warps fight over limited L1/L2, raising the miss rate

5.3 When occupancy matters

Occupancy matters most for kernels that are memory-heavy with little reuse:

• ✅ Elementwise ops (vector add, activations, etc.): almost no reuse; you rely on latency hiding
• ✅ Reductions: each element is read once
• ✅ Simple stencil ops

5.4 When low occupancy is fine

Low occupancy is reasonable when:

• ✅ GEMM / matrix multiply: heavy reuse; registers and shared memory buy compute efficiency
• ✅ FlashAttention: more shared memory cuts HBM traffic
• ✅ Compute-bound kernels: the bottleneck is compute, not memory

5.5 Golden rule

$$\text{performance} = f(\text{Occupancy}, \text{ILP}, \text{data reuse}, \text{memory efficiency}, ...)$$

💡 Tip: Do not lock in a fixed occupancy target (like “must hit 75%”). Measure throughput with a profiler. If lowering occupancy raises data reuse or ILP (Instruction-Level Parallelism), overall performance can improve.

---

6. Tuning strategies and tools

6.1 CUDA Occupancy Calculator

NVIDIA provides an Excel Occupancy Calculator: enter architecture and resource use to get theoretical occupancy and limiters.

You can also use the command line:

# Inspect per-kernel resource use at compile time
nvcc -Xptxas -v my_kernel.cu
# Example output:
# ptxas info: Used 40 registers, 16384 bytes smem, 380 bytes cmem[0]

6.2 Occupancy analysis in Nsight Compute

# Key metrics
sm__warps_active.avg.pct_of_peak_sustained_active   # actual active warp ratio
launch__occupancy                                    # theoretical occupancy
launch__registers_per_thread                         # registers per thread
launch__shared_mem_per_block_allocated               # shared memory per block

Nsight Compute also shows what limits occupancy:

Occupancy Limiters:
  Registers:        50%  ← bottleneck
  Shared Memory:    75%
  Block Size:       100%
  Theoretical:      50%

6.3 Practical tuning flow

graph TD
    A["Measure kernel performance"] --> B["Get occupancy from Nsight Compute"]
    B --> C{"Is occupancy the bottleneck?"}
    C -->|"Memory-bound and Occupancy < 50%"| D["Try raising occupancy"]
    C -->|"Compute-bound or already fast enough"| E["Optimize elsewhere"]
    D --> F{"Limiter?"}
    F -->|"Registers"| G["Fewer temporaries / __launch_bounds__"]
    F -->|"Shared memory"| H["Smaller tiles / dynamic allocation"]
    F -->|"Block size"| I["Adjust block dimensions"]
    G --> J["Check whether performance improved"]
    H --> J
    I --> J
    J --> K{"Did performance improve?"}
    K -->|"Yes"| L["Keep the new config"]
    K -->|"No (e.g. register spilling)"| M["Roll back; accept current occupancy"]

6.4 Block size recommendations

---

7. Case studies

7.1 Case: SGEMM tile-size tradeoff

For a matrix-multiply kernel, tile size directly sets register and shared memory use:

// Tile config comparison
// Config A: small tile
// TILE_M=64, TILE_N=64, TILE_K=8
// Each thread computes 4x4 = 16 outputs → ~48 registers
// Shared memory: (64+64)*8*4 = 4 KB
// Occupancy: ~75%

// Config B: large tile
// TILE_M=128, TILE_N=128, TILE_K=8
// Each thread computes 8x8 = 64 outputs → ~128 registers
// Shared memory: (128+128)*8*4 = 8 KB
// Occupancy: ~25%

// Config B has lower occupancy, but:
// - Each load from Global is reused 8 times (vs 4 for Config A)
// - More work per thread, better ILP
// - Actual performance is usually Config B > Config A

7.2 Case: flexible tuning with dynamic shared memory

// Dynamic shared memory: allocate at launch as needed
__global__ void flexible_kernel(float* data, int N) {
    extern __shared__ float smem[];
    // ...
}

// Pick the best config at runtime from device capability
int smem_size;
if (occupancy_at_16kb > 0.5) {
    smem_size = 16 * 1024;
} else {
    smem_size = 8 * 1024;  // fall back
}

// Configurable shared memory vs L1 split (Volta+)
cudaFuncSetAttribute(
    flexible_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize,
    smem_size
);

flexible_kernel<<<grid, block, smem_size>>>(data, N);

---

📝 Summary

---

🎯 Self-check list

• Can hand-compute occupancy from registers, shared memory, and block size
• Can explain why a 100% occupancy kernel is not necessarily faster than 50%
• Can use cudaOccupancyMaxPotentialBlockSize to pick block size automatically
• Can use __launch_bounds__ to cap registers and understand the side effect (spilling)
• Can explain when register spilling happens and how it hits performance
• Can tell whether a kernel is memory-bound or compute-bound, and set occupancy priority accordingly
• Can read occupancy limiters from an Nsight Compute report
• Can trade tile size vs occupancy sensibly for GEMM-like kernels
• Can use dynamic shared memory and cudaFuncSetAttribute to configure shared memory
• Can pick a suitable block size for different kernel types

---

📚 References

• NVIDIA CUDA C++ Programming Guide - Hardware Multithreading
• NVIDIA CUDA C++ Best Practices Guide - Occupancy
• CUDA Occupancy Calculator
• Achieved Occupancy - NVIDIA Developer Blog
• CUTLASS: Fast Linear Algebra in CUDA C++ - GitHub