Why Is My GPU Idle? A CUDA Bottleneck Checklist for Engineers
Low GPU utilization is a symptom report, not a diagnosis. This guide maps the five root causes to the Nsight Compute counters that reveal them — and shows you one concrete fix to validate the loop.
§1 — The Symptom
nvidia-smi reports 23% GPU utilization. You have tried increasing batch size. You have tried pinning more DataLoader workers. You have enabled AMP. The number barely moves.
The problem is that GPU utilization is a utilization rate — what fraction of time the GPU was doing something. It cannot tell you why it is idle. That answer requires a profiler and a mental model of the CUDA execution pipeline.
Six different root causes produce identical utilization readings from the outside:
Six causes that look the same from nvidia-smi
- CPU feeding data too slowly (input starvation)
- Host-to-device transfer dominating the timeline
- Too many small kernel launches overwhelming the scheduler
- Memory bandwidth saturated before compute peaks
- Low occupancy leaving SM warp slots empty
- Synchronization barriers stalling the entire device
Each one has a different fix. Applying the wrong one wastes a sprint. The right approach is to profile first and let the counters point you to the cause.
§2 — The CUDA Performance Loop
NVIDIA’s CUDA Best Practices Guide frames the optimization loop as APOD. The critical point is that it is a loop, not a waterfall — after Deploy you return to Assess and peel off the next constraint.
Assess
Profile the workload first. Measure where time actually goes before touching a line of code.
Parallelize
Port the hot path to CUDA kernels. Most of the gain usually comes from parallelism alone.
Optimize
Address the specific bottleneck the profiler identified — memory, occupancy, or launch overhead.
Deploy
Measure again. Confirm the gain is real, then return to Assess for the next constraint.
§3 — Five Common CUDA Bottlenecks
These five bottlenecks account for the vast majority of GPU underutilization in real training and inference workloads. They are not mutually exclusive — a single kernel can simultaneously suffer from uncoalesced memory access and low occupancy.
Host-device transfer overhead
Every pageable cudaMemcpy blocks the CPU until the transfer completes and prevents the GPU from running compute in that stream. Workloads that constantly shuttle data between RAM and VRAM make the PCIe bus the actual bottleneck, not the GPU.
H2D/D2H copy time dominates; compute gaps visible in timeline
Uncoalesced global memory access
A warp of 32 threads reading 32 contiguous floats generates one memory transaction. The same warp reading every 32nd float generates 32 transactions. That 32× amplification destroys DRAM throughput even when arithmetic intensity is low.
Global load sectors/request > 4; DRAM throughput high but compute low
Low occupancy
The SM scheduler hides memory latency by switching between resident warps. If too few warps are resident — because kernels use too many registers, claim too much shared memory, or launch with a tiny block size — the scheduler stalls with nothing to run.
Achieved occupancy < 50%; warp stall reasons show 'no active warps'
Synchronization stalls
__syncthreads() forces all threads in a block to reach the barrier before any can continue. When threads diverge, fast threads idle at the barrier. The entire GPU can appear underutilized even though every SM reports 'active.'
Dominant warp stall reason = sync; __syncthreads() density high vs. compute
Launch overhead / many tiny kernels
Each kernel launch incurs ~5–20 µs of CPU-side driver overhead. A model with thousands of small element-wise ops — common in eager-mode PyTorch before graph capture — spends more wall time in launch overhead than in actual compute.
Kernel median duration < 10 µs; nvidia-smi shows brief bursts with long idle gaps
§4 — What to Look for in Nsight Compute
Nsight Compute exposes hundreds of hardware counters. Most runs you only need seven. These map directly to the five bottleneck categories above and give you a threshold to anchor your diagnosis. Reference: Nsight Compute Profiling Guide.
| Counter | What it measures | Healthy | Problem signal |
|---|---|---|---|
| DRAM Throughput % | GB/s consumed vs. peak memory bandwidth | < 70% for compute-bound kernels | ≥ 80% → kernel is memory-bandwidth-bound |
| L1 Hit Rate % | Requests served from L1 vs. forwarded to L2/DRAM | > 60% | < 40% → cache thrashing or large strided access |
| L2 Hit Rate % | Requests served from L2 vs. forwarded to DRAM | > 65% | < 50% → working set exceeds L2 capacity |
| Global Load Sectors/Request | DRAM transactions generated per global load instruction | ≤ 4 (1 cache line per warp) | > 4 → uncoalesced or strided access pattern |
| Achieved Occupancy % | Active warps vs. theoretical warp capacity | > 50% (higher is not always better) | < 30% → register/shared-memory pressure or small block size |
| Issue Slot Utilization % | Instructions issued per available scheduler cycle | > 60% | Low → scheduler has warps but nothing is ready to issue |
| Warp Stall Reasons | Why warps are not issuing instructions | No single stall > 10% | Dominant stall category directly names the root cause |
The Warp Stall Reasons counter is especially useful: its dominant category directly names the bottleneck. A kernel with stall_memory_throttle as the largest stall is memory-bandwidth-bound. A kernel with stall_sync dominant has too many synchronization barriers. No manual threshold-checking needed — the profiler tells you.
§5 — One Concrete Fix: Pinned Memory + Async Copies
Host-device transfer overhead is one of the most common and most fixable bottlenecks. The standard cudaMemcpy with pageable memory is synchronous — it blocks the CPU thread and prevents the GPU from doing compute in that stream simultaneously. Switching to cudaMallocHost (pinned memory) and cudaMemcpyAsync with a stream lets the GPU begin executing the kernel as soon as the data is available, while the CPU is free to prepare the next batch. Reference: How to Overlap Data Transfers in CUDA C/C++.
Before — pageable memory, blocking copy
// ❌ Pageable memory — CPU blocks until copy completes
float* h_data = (float*)malloc(N * sizeof(float));
fill_data(h_data, N);
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
myKernel<<<grid, block>>>(d_data, N);
cudaDeviceSynchronize();
free(h_data);After — pinned memory, async copy, streams
// ✅ Pinned memory + async copy — overlaps transfer with compute
float* h_data;
cudaMallocHost(&h_data, N * sizeof(float)); // page-locked allocation
fill_data(h_data, N);
cudaStream_t stream;
cudaStreamCreate(&stream);
// Transfer starts immediately; CPU is NOT blocked
cudaMemcpyAsync(d_data, h_data, N * sizeof(float),
cudaMemcpyHostToDevice, stream);
// Kernel executes in same stream, so it waits for the copy —
// but a *second* stream can overlap with both
myKernel<<<grid, block, 0, stream>>>(d_data, N);
// Pre-process the NEXT batch on CPU while GPU works
fill_data(h_next, N);
cudaStreamSynchronize(stream);
cudaFreeHost(h_data);After applying the fix, re-run the profiler and check the same counters. DRAM throughput should drop (less time copying), kernel occupancy should hold steady, and wall-clock throughput should rise. Here is what the frx profile output looks like on a memory-bandwidth-bound kernel — before the layout fix but after pinned memory removes the transfer stall:
$ frx profile --ncu report.csv
VERDICT memory_bandwidth_bound (score 0.87)
MEASURED METRICS
DRAM Throughput 91.3 % [! above 70 % threshold]
L2 Hit Rate 54.1 % [ok]
Global Load Sectors/Req 2.1 [ok]
Achieved Occupancy 72.4 % [ok]
BOTTLENECKS DETECTED
████████████████████ memory_bandwidth_bound 0.87
███░░░░░░░░░░░░░░░░░ l2_cache_thrashing 0.31
RECOMMENDATIONS
[HIGH] Improve memory access coalescing
Why: DRAM throughput is 91 % of peak — kernel is bandwidth-bound
Actions: 1. Ensure adjacent threads access adjacent memory addresses
2. Use shared-memory tiling to reuse loaded data
3. Consider AoS → SoA layout change
Validate: rerun with --preset memory; target DRAM throughput < 70 %
Estimated speedup: 20–60 %
NEXT STEPS
frx profile --ncu report.csv --preset memory§6 — Where Fournex Fits
The workflow above — profile, read counters, map to root cause, pick a fix, validate — is the right mental model. But in production, with dozens of kernels across a multi-GPU training job, the “read counters” step quickly becomes a wall of CSV rows with no obvious signal.
Fournex automates the mapping step. You run frx profile against your Nsight Compute export, and it labels the bottleneck, ranks the evidence, and outputs the next optimization move — with the estimated speedup and validation command included.
Fournex labels the bottleneck,
ranks the evidence, ships the fix.
- Labels the bottleneck from raw NCU counters — no manual threshold-checking
- Ranks fixes by estimated speedup, evidence confidence, and implementation effort
- Generates a ready-to-paste LLM brief for every kernel that needs attention