Profiling CUDA: Occupancy, Memory Coalescing, and Nsight

A CUDA kernel can be perfectly correct and still hit ~3% of peak FLOPS. Without a profiler you'll guess wrong about why. The three usual suspects:

• Memory-bound vs compute-bound. Most kernels are memory-bound; the GPU is idle waiting for VRAM.
• Low occupancy. Not enough warps resident on each SM to hide memory latency.
• Warp divergence. Branches splitting warps, halving (or worse) effective throughput.

The NVIDIA toolchain gives you three windows into this: nvcc --ptxas-options=-v at compile time, nvprof/nsys for high-level timing, and Nsight Compute for per-kernel deep dives. Learn the first and third; the second is mostly bookkeeping.

An SM has a fixed register file (e.g. 65,536 regs) and shared memory budget (e.g. 100 KB), and a hardware limit on resident warps (e.g. 64 warps = 2048 threads).

Occupancy is the ratio of active warps per SM to that hardware max. Higher occupancy means the warp scheduler has more candidates to swap in whenever one is stalled on memory — that's how the GPU hides latency.

Things that cap occupancy:

• Registers per thread. Using 64 regs/thread on a 65k-reg SM caps you at 1024 threads = 50% occupancy.
• Shared memory per block. Asking for 48 KB per block when an SM has 100 KB means only 2 blocks resident.
• Block size. Tiny blocks (e.g. 32 threads) leave most of the SM idle.

Aim for ≥50% achieved occupancy on memory-bound kernels. Compute-bound kernels can be fast at lower occupancy.

Before you launch anything, ask the compiler what it produced:

nvcc -O3 --ptxas-options=-v kernel.cu -o kernel

You'll get a line per kernel like:

ptxas info    : Used 38 registers, 8192 bytes smem, 0 bytes lmem, 360 bytes cmem[0]

What to read:

• registers: 38/thread × 256 threads/block = 9728 regs/block. SM has 65,536 → up to 6 blocks resident. Healthy.
• smem: shared memory per block. Crosscheck against SM budget.
• lmem: local memory spills. Should be 0. Non-zero means the compiler couldn't keep something in registers and moved it to global VRAM. Painful and almost always fixable.
• cmem: constant memory usage. Usually fine.

If you see lmem > 0, look for large local arrays with dynamic indices or excessive register pressure. Either reduce register usage with __launch_bounds__ or restructure the data.

The single biggest performance lever on memory-bound kernels is coalescing. A warp of 32 threads issuing 32 loads should produce one 128-byte transaction, not 32.

// Coalesced: 32 threads -> 32 contiguous floats -> 1 transaction
float v = a[blockIdx.x * blockDim.x + threadIdx.x];

// Strided: 32 threads -> 32 floats with stride 16 -> 32 transactions
float v = a[(blockIdx.x * blockDim.x + threadIdx.x) * 16];

For 2D data the lesson is: threadIdx.x should index the fastest-varying dimension. If your matrix is row-major and you want to iterate column-wise per warp, you've built strided access. The fix is usually swapping threadIdx.x/threadIdx.y roles or transposing the data once into shared memory.

Same compute, very different memory traffic.

ncu --set full kernel (or the GUI) gives you dozens of metrics. The triage list for a slow kernel:

1. Achieved Occupancy — actual active warps / max. Below 25%? Cut register/shared usage or grow block count.
2. Memory Throughput (% of peak bandwidth) and Compute Throughput (% of peak FLOPs). Whichever is higher tells you the bottleneck. If both are <30%, you're stalled on something else.
3. Warp Stall Reasons — why the warps were waiting. Long Scoreboard = global memory latency. Short Scoreboard = shared memory or register dependency. Barrier = __syncthreads() waits. MIO Throttle = uncoalesced access.

A classic profile: 80% memory throughput, 5% compute throughput, dominant stall = Long Scoreboard. Diagnosis: memory-bound, coalescing is OK, you need more reuse — try shared memory tiling.

When threads in a warp take different paths, the hardware runs every path serially with the inactive threads masked off. Nsight reports it as Branch Efficiency (lower = worse) and as Branch stalls.

The common offenders:

// Bad: half the warp goes one way, half the other
if (threadIdx.x < 16) doA(); else doB();

// Bad: data-dependent branch with no warp coherence
if (a[i] > 0) doA(); else doB();

// OK: branch correlates with blockIdx — whole warp agrees
if (blockIdx.y == 0) doEdgeCase(); else doInner();

Fixes: hoist the branch above the warp (split into two kernels), use predicated arithmetic instead of branches (x = cond ? a : b), or sort data so warps see coherent values.

When a kernel is slow, attack it in this order — earlier fixes are cheaper and have bigger upside:

1. Fix correctness and bounds first. A profiler reading on a broken kernel is misleading.
2. Coalesce global memory. Biggest single win on most kernels.
3. Eliminate local-memory spills. Check ptxas -v for lmem.
4. Add shared memory reuse if the same data is touched multiple times.
5. Tune block size for occupancy — usually 128, 256, or 512.
6. Reduce warp divergence by restructuring branches.
7. Overlap compute and transfer with streams and cudaMemcpyAsync.
8. Last-mile: bank conflicts, instruction-level parallelism, tensor cores.

Most kernels live or die on steps 2-3. Steps 7-8 matter for the last 2x once you're already at 60%+ of peak.

Suppose Nsight reports:

• Achieved Occupancy: 18%
• Memory Throughput: 22%
• Compute Throughput: 4%
• Top stall reason: Long Scoreboard (40%)
• ptxas -v: 96 regs/thread, 0 bytes lmem

Reading: occupancy is killed by register pressure (96 × 256 = 24,576 regs/block; SM holds at most 2 blocks). The kernel is memory-stalled (Long Scoreboard), but the GPU has no spare warps to swap in.

Fix path:

1. Cap registers with __launch_bounds__(256, 4) to force the compiler to spill carefully and target 4 blocks/SM. Re-measure.
2. If memory throughput stays low even with healthier occupancy, the kernel's access pattern is genuinely sparse — consider rewriting with shared-memory tiling or texture caching.

This is the loop: measure, hypothesise, change one thing, measure again.