Occupancy & Resource Allocation

Occupancy determines how many warps can run simultaneously on an SM, which determines how effectively the GPU hides memory latency. This chapter covers the resource model, worked calculations, and the counterintuitive truth that maximum occupancy isn’t always maximum performance.

---

1. What Occupancy Means

$$\text{Occupancy} = \frac{\text{Active warps per SM}}{\text{Maximum warps per SM}}$$

H100: maximum 64 warps (2048 threads) per SM. If your kernel can only run 32 warps per SM, occupancy is 50%.

Why it matters: More active warps → more warps available for the scheduler to switch between → better latency hiding when warps stall on memory accesses.

---

2. SM Resource Budget

Each SM has a fixed pool of resources. Blocks (and their warps) are assigned to an SM until any resource is exhausted:

Resource	H100 per SM	A100 per SM	Shared across
Max threads	2048	2048	All blocks on the SM
Max warps	64	64	All blocks
Max blocks	32	32	—
Registers	65,536 (256 KB)	65,536 (256 KB)	All threads on the SM
Shared memory	228 KB	164 KB	All blocks on the SM
Max registers/thread	255	255	—
Max threads/block	1024	1024	—

Resource Allocation is Per-Block

When a block is assigned to an SM:

• Its registers are deducted from the SM register pool (per-thread × threads-in-block)
• Its shared memory is deducted from the SM shared memory pool
• Its thread slots are deducted from the SM thread limit

The SM keeps accepting blocks until any one resource would be exceeded.

---

3. The Three Limiters

Limiter 1: Registers

Registers per SM: 65,536
Registers per thread: 48 (example)
Threads per block: 256

Registers consumed per block: 48 × 256 = 12,288
Max blocks by registers: 65,536 / 12,288 = 5 blocks
Active warps: 5 × 8 = 40 → Occupancy: 40/64 = 62.5%

If registers per thread = 96:

Registers per block: 96 × 256 = 24,576
Max blocks: 65,536 / 24,576 = 2 blocks
Active warps: 2 × 8 = 16 → Occupancy: 25% ← register-limited!

Limiter 2: Shared Memory

Shared memory per SM: 228 KB
Shared memory per block: 48 KB

Max blocks by shared memory: 228 / 48 = 4 blocks
Active warps: 4 × 8 = 32 → Occupancy: 50%

Limiter 3: Block Size

Max threads per SM: 2048
Threads per block: 512

Max blocks: 2048 / 512 = 4 blocks
Active warps: 4 × 16 = 64 → Occupancy: 100%

But if threads per block = 1024:
Max blocks: 2048 / 1024 = 2 blocks
Active warps: 2 × 32 = 64 → Still 100% (larger blocks, fewer of them)

---

4. Worked Example: Finding the Bottleneck

A kernel uses 64 registers/thread, 16 KB shared memory/block, 256 threads/block:

Resource	Limit	Blocks allowed
Registers	65,536 / (64 × 256) = 4	4
Shared memory	228 KB / 16 KB = 14	14
Threads	2048 / 256 = 8	8
Block limit	32	32
Actual	min(4, 14, 8, 32)	4 (register-limited)

Occupancy: 4 blocks × 8 warps = 32 active warps / 64 max = 50%

To improve: reduce register usage. If we get to 48 registers/thread:

• 65,536 / (48 × 256) = 5 blocks → 40 warps → 62.5%

---

5. Using the Occupancy API

Compile-Time Analysis

nvcc --ptxas-options=-v kernel.cu
# Output: Used 48 registers, 16384 bytes smem, 0 bytes cmem

Runtime Optimal Block Size

int minGridSize, bestBlockSize;
cudaOccupancyMaxPotentialBlockSize(
    &minGridSize, &bestBlockSize,
    myKernel,
    0,    // dynamic shared mem (0 if static)
    0     // block size limit (0 = no limit)
);
printf("Optimal block size: %d, min grid: %d\n", bestBlockSize, minGridSize);

Runtime Occupancy Query

int maxActiveBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &maxActiveBlocks, myKernel, blockSize, sharedMemSize
);
float occupancy = (float)(maxActiveBlocks * blockSize / 32) / 64.0f;
printf("Occupancy: %.1f%%\n", occupancy * 100);

---

6. When High Occupancy Doesn’t Help

Compute-Bound Kernels

If a kernel has high arithmetic intensity (large GEMM), every warp is busy computing — there’s nothing to hide. More warps just compete for the same compute units.

Example: A GEMM kernel at 25% occupancy achieves 90% of peak TFLOPS because it’s compute-bound. Increasing to 50% occupancy doesn’t improve throughput — it just adds register pressure.

Diminishing Returns

Going from 25% → 50% occupancy can give 2× speedup (for memory-bound kernels). Going from 50% → 100% typically gives <20% improvement — once you have ~30 active warps, the scheduler has enough to keep the pipeline full.

The Tradeoff

Higher occupancy often requires:

• Fewer registers per thread → more register spilling → slower
• Less shared memory per block → less data reuse → slower

The optimal point: enough occupancy to hide latency, but not so much that you sacrifice per-thread efficiency.

Rule of thumb: Profile at 50% and 100% occupancy. If the difference is <10%, keep the lower occupancy (fewer registers = faster per-thread execution).

---

7. Occupancy Optimization Strategies

Strategy	Effect on occupancy	Tradeoff
Reduce registers (-maxrregcount=N)	↑ More threads fit	May cause spilling
Reduce shared memory	↑ More blocks fit	Less data reuse
Reduce block size	↑ or ↓ depending	Fewer warps per block
Use __launch_bounds__	Compiler optimizes for target	Limits flexibility

// Tell compiler to optimize for 256 threads, at least 4 blocks per SM
__global__ __launch_bounds__(256, 4)
void myKernel(...) { ... }

---

Self-Check

• Define occupancy and explain why it matters for latency hiding
• Calculate occupancy given register count, shared memory, and block size
• Identify which resource (registers, shared memory, threads) is the bottleneck
• Use cudaOccupancyMaxPotentialBlockSize to find the optimal block size
• Explain why 100% occupancy can be slower than 50% for compute-bound kernels
• Use __launch_bounds__ to help the compiler optimize

---

References

• CUDA Occupancy Calculator
• CUDA Best Practices — Occupancy
• Achieved Occupancy Metric Guide