---
name: gpu-coalescing-and-occupancy
description: The two knobs that dominate GPU perf — memory coalescing + occupancy targets across CUDA (Blackwell), Metal (Apple 7+), ROCm (RDNA3/CDNA3).
when_to_use: You're a kernel_author or bench_engineer profiling a GPU kernel. GPU team pins.
tags: [gpu, cuda, metal, rocm, performance]
---

# Coalescing + occupancy

Two dominant knobs on ANY modern GPU. Get them right first; then worry about smaller wins.

## Coalescing

**Rule**: consecutive threads in a warp/wavefront should read consecutive 32/64/128-bit words.

- **CUDA (Blackwell, SM_100+)** — warp = 32 threads. Threads 0..31 accessing `arr[tid]` = 1 memory transaction. `arr[tid * stride]` for stride > 1 = multiple transactions. Violation cost: 4×–32× slowdown depending on stride.
- **Metal (Apple 7+, i.e., A17 / M3+)** — SIMD-group = 32 threads, same principle. `metal::simd_shuffle` for cross-lane comms without shared memory.
- **ROCm (RDNA3, CDNA3)** — wavefront = 32 (RDNA) or 64 (CDNA). Same rule; note the 64-wide wavefronts on CDNA change your indexing math.

Prove coalescing in a comment on the kernel:
```cpp
// Coalesced: thread tid reads global[tid], stride-1.
// Warp-level transaction: 1× 128-byte load per warp.
float x = global_in[gid];
```

## Occupancy

Rough targets:
- **Memory-bound kernel**: aim for **≥ 50% occupancy** to hide DRAM latency.
- **Compute-bound kernel with high ILP**: 25%–50% is often fine; more threads compete for registers + shared memory.

Occupancy is bounded by whichever hits first:
- Registers per thread × threads per block ≤ register file per SM.
- Shared / threadgroup memory per block ≤ per-SM budget.
- Threads per block ≤ max (typically 1024).

Tools:
- **Nsight Compute** (CUDA): `--section=Occupancy` prints the bottleneck.
- **Xcode GPU Frame Capture** (Metal): occupancy chart in the performance report.
- **rocprof** (ROCm): `--stats` for kernel occupancy.

## Memory hierarchy strategy

Order-of-magnitude comparison — always prefer the fastest tier your data footprint allows:

| Tier | CUDA name | Latency | Notes |
|---|---|---|---|
| Register | reg | ~1 cycle | Per-thread; too many spills to L1 |
| Shared | shared / threadgroup / LDS | ~30 cycles | Per-block. CUDA: 48–100 KB. Metal: 32 KB (Apple7+). ROCm: 64 KB LDS |
| L1 / texture | L1 tex | ~200 cycles | Cache; some HW has separate paths |
| L2 | L2 | ~500 cycles | Shared across SMs |
| Global | DRAM / HBM | 400–800 cycles | Coalescing matters MOST here |

## Divergence

- Warp-uniform control flow keeps all lanes active.
- `if (tid % 2)` splits the warp — half the lanes idle per branch.
- On CUDA Volta+ (7.0+), Independent Thread Scheduling makes divergence CORRECT under all conditions but doesn't make it FAST.

## The one bottleneck rule

After profiling, report ONE bottleneck to attack next. Reporting five doesn't help — you'll fix one and re-profile anyway.

Format:
> `KERNEL foo achieves 340 GB/s of 3.35 TB/s HBM = 10% BW-bound. Occupancy 62%. Bottleneck: stride-2 loads from arr — refactor tile shape to coalesce.`