Understanding GPU memory bottlenecks
Most "slow" kernels aren't compute-bound — they're wasting bandwidth. A tour of the seven ways VRAM quietly costs you money.
When a kernel is slow, the instinct is to blame compute. Usually it’s memory. Modern GPUs have enormous arithmetic throughput and comparatively scarce bandwidth, so most real kernels spend their time waiting on data, not crunching it. Here are the seven memory inefficiencies Kernalin’s analyzer looks for — and why each one costs you.
1. Fragmentation
The caching allocator reserves VRAM in blocks. Over time, blocks get split into pieces too small to reuse, and you end up with gigabytes “reserved” but unusable. The fix is usually expandable_segments:True, or max_split_size_mb when it can’t be used.
2. Uncoalesced access
When 32 threads in a warp read scattered addresses, the hardware fetches whole 32-byte sectors and throws most of each away. The signal is low bytes-per-sector efficiency. The fix is a layout change so consecutive threads touch consecutive addresses — and it’s width-aware, because a perfectly coalesced float4 load legitimately touches more sectors.
3. Excessive allocations
Allocating fresh tensors every iteration thrashes the allocator and can trigger retries and OOMs. Reuse buffers, pass out=, and capture the step in a CUDA graph.
4. Poor cache locality
Low L1/L2 hit rates mean you’re re-reading from DRAM what you could have kept in cache. Tiling and fusion keep the working set resident.
5. DRAM bandwidth saturation
If you’re pinned at ~85% of peak DRAM throughput with poor reuse, you’re streaming — and the fix is to move fewer bytes (fusion, lower precision) rather than chase occupancy.
6. Shared-memory bank conflicts
When threads hit the same shared-memory bank, accesses serialize into replays. Padding arrays by one column or swizzling indices usually clears it.
7. Register spilling
When a kernel needs more registers than it has, they spill to local memory — which is just DRAM with a friendly name. Any local-memory traffic on a compute kernel is a red flag; cut register pressure with smaller tiles or __launch_bounds__.
Diagnosis, not vibes
The point of naming all seven is that “it’s memory-bound” isn’t actionable — but “30% of reserved VRAM is stranded in split blocks, set expandable_segments:True” is. Kernalin reports each dimension with the number that triggered it and a concrete, prioritized recommendation.
This is sample content published to illustrate Kernalin’s engineering perspective.