Occupancy & Launch Configuration

Occupancy = active warps per SM ÷ maximum warps per SM (e.g., 64 on A100/H100).

20.1 Occupancy

Occupancy = active warps per SM ÷ maximum warps per SM (e.g., 64 on A100/H100). Limited by:

Registers per thread (64K registers per SM shared by all resident threads).

Shared memory per block.

Threads per block and max blocks per SM (16–32).

Example: 256 threads/block × 64 registers/thread = 16K registers/block → 4 blocks/SM → 1024 threads = 32 warps → 50% occupancy (of 64).

Higher occupancy helps hide latency, but isn't the goal itself — some of the fastest kernels run at low occupancy with high ILP and register reuse.

20.2 Occupancy API

Nsight Compute shows achieved vs theoretical occupancy and limiting factors.

int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, myKernel, 0, 0);
int gridSize = (n + blockSize - 1) / blockSize;
myKernel<<<gridSize, blockSize>>>(...);

int numBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, myKernel, blockSize, sharedBytes);

20.3 Controlling register usage

Too few registers → spills to local memory (check -Xptxas -v for "spill stores").

__global__ void __launch_bounds__(256, 2) kernel(...);   // max 256 threads/block, aim for ≥2 blocks/SM
// or nvcc -maxrregcount=64 (global)

20.4 Launch configuration guidelines

Threads per block: multiple of 32; start with 128–256.

Grid: at least enough blocks to fill all SMs several times, or a grid-stride loop with k × numSMs blocks.

Tail effects: a grid that doesn't divide evenly over SMs wastes the last "wave".

Measure — don't guess.

🧪 Part II Exercises

Measure effective bandwidth of a coalesced copy vs a strided copy (stride 2, 4, 8, 32).

Implement naive, shared-memory, and padded shared-memory matrix transpose; compare bandwidth.

Implement tiled matrix multiplication for arbitrary N; compare against the naive version and cuBLAS SGEMM.

Implement a sum reduction in four stages of optimization (interleaved → sequential → warp shuffle → grid-stride + atomics); report GB/s for each.

Write an inclusive scan for 10⁸ integers (multi-block) and verify against std::inclusive_scan.

Implement a 256-bin histogram with global atomics vs shared-memory privatization; compare speed on random and skewed data.

Write a 2-D Gaussian blur using constant memory for the kernel and shared-memory tiles with halos.

Overlap transfers and computation with 1, 2, 4, 8 streams using pinned memory; visualize with Nsight Systems.

Port vecAdd to unified memory; measure with and without cudaMemPrefetchAsync.

Use the occupancy API to choose block sizes for three kernels; compare with your manual choice.