CUDA Profiling with Nsight: The Complete 2026 Guide to GPU Performance Engineering
1. Why GPU Profiling Is Non-Negotiable in 2026 {#why-profiling}
You optimized your CUDA kernel from 10 ms → 8 ms. Great. But what if the real problem isn't the kernel at all?
1Kernel execution 8 ms 2Memory transfer 20 ms 3CPU overhead 15 ms 4Synchronization 10 ms
In modern AI systems — from LLM inference to computer vision training — naive optimization is expensive guesswork. Without profiling, you're tuning in the dark.
The difference between a CUDA programmer and a GPU performance engineer is the profiler.
This guide covers the two most important NVIDIA profiling tools in 2026:
- Nsight Systems (
nsys) — System-wide timeline view - Nsight Compute (
ncu) — Kernel-level hardware counter analysis
2. Nsight Systems vs Nsight Compute: The Mental Model {#nsight-vs-nsight}
| Question | Use This Tool |
|---|---|
| "Why is my app slow overall?" | Nsight Systems |
| "Why is THIS specific kernel slow?" | Nsight Compute |
| "Are my streams overlapping?" | Nsight Systems |
| "Is my kernel memory-bound or compute-bound?" | Nsight Compute |
| "Where are the CPU-GPU synchronization gaps?" | Nsight Systems |
| "Are my Tensor Cores actually firing?" | Nsight Compute |
Remember:
1Nsight Systems → WHOLE APPLICATION → Timeline 2Nsight Compute → ONE KERNEL → Detailed Metrics
3. The Golden Profiling Loop {#golden-loop}
Every GPU performance engineer follows this loop. Memorize it:
1┌──────────────────┐ 2│ Run Application │ 3└────────┬─────────┘ 4 ↓ 5┌──────────────────┐ 6│ Nsight Systems │ ← Find the hot kernel 7└────────┬─────────┘ 8 ↓ 9┌──────────────────┐ 10│ Nsight Compute │ ← Find the bottleneck 11└────────┬─────────┘ 12 ↓ 13┌──────────────────┐ 14│ Change Kernel │ ← Optimize based on data 15└────────┬─────────┘ 16 ↓ 17┌──────────────────┐ 18│ Benchmark Again │ ← Validate improvement 19└────────┬─────────┘ 20 ↓ 21 Repeat
Bad workflow: "I heard shared memory is faster → Rewrite kernel."
Good workflow: Profile → Identify memory bottleneck → Analyze access pattern → Optimize → Measure.
4. Nsight Systems: System-Level Timeline Analysis {#nsight-systems}
Basic Command
1nsys profile -o report ./my_program
For AI/ML workloads, capture CUDA, NVTX, and OS runtime:
1nsys profile \ 2 --trace=cuda,nvtx,osrt,cudnn \ 3 -o sys_profile \ 4 python train_model.py
What the Timeline Reveals
Conceptually, Nsight Systems shows:
1Time → 2──────────────────────────────────── 3 4CPU: 5Launch A ── Launch B ── Wait ────── 6 7GPU: 8H2D █████ 9 Kernel A █████████ 10 Kernel B ██████ 11 D2H █████
What to Look For
| Pattern | Meaning | Action |
|---|---|---|
| Large GPU idle gaps | CPU bottleneck, sync issues, or launch overhead | Check CPU thread activity |
| Serialized kernels | Missing stream concurrency or false dependencies | Use multiple CUDA streams |
| Back-to-back memory copies | Excessive H2D/D2H traffic | Use pinned memory, batch transfers |
| CPU launch too slow | GPU starved for work | Reduce launch overhead, use CUDA Graphs |
| No stream overlap | Streams not actually concurrent | Verify async copies & independent streams |
Pro Tip: Remote Cloud Profiling
1# Capture headlessly on server 2nsys profile --trace=cuda,nvtx,osrt -o profile.nsys-rep python inference.py 3 4# Download and analyze locally 5scp user@host:~/profile.nsys-rep ./ 6nsys-ui profile.nsys-rep
5. Nsight Compute: Kernel-Level Deep Dive {#nsight-compute}
Basic Command
1ncu -o kernel_report ./my_program
Scoped Profiling (Essential for AI Workloads)
Don't profile everything. Target the hot kernel:
1ncu \ 2 --replay-mode kernel \ 3 --set full \ 4 --kernel-name flash_attn_varlen_fwd \ 5 --launch-skip 10 \ 6 --launch-count 5 \ 7 -o attn_profile \ 8 python inference.py
What Nsight Compute Analyzes
1┌───────────────────────────┐ 2│ Kernel Performance │ 3├───────────────────────────┤ 4│ GPU utilization │ 5│ SM throughput │ 6│ Memory throughput │ 7│ Warp execution │ 8│ Occupancy │ 9│ Registers │ 10│ Shared memory │ 11│ Cache behavior │ 12│ Instructions │ 13│ Tensor Cores │ 14│ Stall reasons │ 15└───────────────────────────┘
6. Memory-Bound vs Compute-Bound: The Roofline Model {#roofline}
The Roofline model is the most important conceptual tool in GPU profiling.
Performance (GFLOP/s)
↑
│ ═══════════════════ Compute Ceiling
│ /
│ /
│ /
│ /
│ /
│ /
└────────────────────────────────→
Arithmetic Intensity (FLOPs/Byte)
Arithmetic Intensity = FLOPs / Bytes of DRAM traffic
H100 SXM5 Reference (2026)
- HBM3 Bandwidth: 3.35 TB/s
- BF16 Compute: 989 TFLOP/s
- Ridge Point: ~295 FLOP/byte
Where AI Kernels Land
| Kernel Type | Arithmetic Intensity | Classification | Fix If Slow |
|---|---|---|---|
| Large batch GEMM | High (~1000+) | Compute-bound | Tensor Cores, tiling |
| Attention decode (batch=1) | Very low (~4-8) | Memory-bound | FlashAttention, increase batch |
| LayerNorm, RoPE, elementwise | Extremely low (<1) | Memory-bound | Kernel fusion |
| FP8 GEMM (medium batch) | ~200-300 | Near ridge | Optimize both memory & compute |
7. Key Metrics Every CUDA Engineer Must Know {#key-metrics}
1. Kernel Duration
1Kernel A = 2 μs 2Kernel B = 500 μs ← Optimize THIS first 3Kernel C = 4 μs
Rule: Optimize the biggest contributors first.
2. GPU Utilization
- 40% utilization → GPU isn't kept busy. Check CPU bottlenecks, sync, or small kernels.
3. SM Utilization
- High SM util ≠ optimal performance. The SMs might be busy stalling, not computing.
4. Memory Throughput
- High DRAM throughput + low compute = memory-bound kernel
- Focus on: coalescing, shared memory, cache, data reuse
5. Occupancy
1Occupancy = Active Warps / Maximum Possible Warps
⚠️ Critical Warning:
50% occupancy can outperform 100% occupancy if the kernel efficiently uses available resources. Occupancy is a diagnostic metric, not the final goal.
6. Warp Execution & Divergence
1if (threadIdx.x % 2) { A(); } else { B(); }
This causes warp divergence — measurable in Nsight Compute.
7. Register Pressure
1Registers/thread = 120 → High pressure → Lower occupancy → Potential spilling
Register spilling → Local memory traffic → Hidden performance killer.
8. Shared Memory Bank Conflicts
1__shared__ float tile[32][32]; // May cause bank conflicts 2__shared__ float tile[32][33]; // Padded — often fixes conflicts
9. Cache Behavior
- L1/L2 hit rates reveal reuse efficiency
- Poor reuse = repeated DRAM fetches
10. Tensor Core Utilization
Expected:
1Tensor Core ███████████████
Actual (sometimes):
1Tensor Core ███
If low, check: datatype, dimensions, memory bottleneck, kernel selection.
8. Profiling AI Workloads: GEMM, Attention & LayerNorm {#ai-workloads}
Example 1: GEMM Profiling
1C = A @ B # FP16, large matrices
Profile checklist:
- Check datatype (FP16/BF16/FP8?)
- Check matrix dimensions
- Check memory throughput
- Check Tensor Core utilization
- Check occupancy
- Compare against cuBLAS/cuBLASLt
If your custom kernel is slower than cuBLAS, the library is already doing sophisticated tiling and scheduling. Use it.
Example 2: Attention Kernel
1QKᵀ → store → softmax → store → × V
Profiling insight: Large memory traffic between ops.
Fix: Fuse operations (FlashAttention) to keep intermediate values in SRAM instead of HBM.
Example 3: LayerNorm
Profile shows: High DRAM traffic, low compute utilization.
Classification: Memory-bound.
Fix: Fuse with adjacent ops (scale + bias) to reduce kernel launches and HBM trips.
9. Practical Profiling Commands & Workflows {#commands}
Complete Nsight Systems Workflow
1# 1. Profile 2nsys profile \ 3 --trace=cuda,nvtx,osrt,cudnn \ 4 --gpu-metrics-device=all \ 5 -o app_profile \ 6 python train.py 7 8# 2. Quick stats in terminal 9nsys stats app_profile.nsys-rep 10 11# 3. Deep visual analysis (local GUI) 12nsys-ui app_profile.nsys-rep
Complete Nsight Compute Workflow
1# 1. Default set first (fast, checks roofline basics) 2ncu --set default -o quick_profile ./my_program 3 4# 2. Full analysis on hot kernel 5ncu \ 6 --replay-mode kernel \ 7 --set full \ 8 --kernel-name my_kernel \ 9 --launch-skip 5 \ 10 --launch-count 3 \ 11 -o deep_profile \ 12 ./my_program 13 14# 3. Open in GUI 15ncu-ui deep_profile.ncu-rep
Docker Profiling Setup (2026)
1FROM nvcr.io/nvidia/pytorch:24.10-py3 2RUN pip install HolisticTraceAnalysis nvidia-pytool-report triton==3.0.0
Run with hardware counter access:
1docker run --rm --gpus all --cap-add SYS_ADMIN \ 2 -v $(pwd):/workspace \ 3 profiling-image \ 4 ncu --replay-mode kernel --set full \ 5 -o /workspace/profile.ncu-rep \ 6 python /workspace/inference.py
10. Common Bottlenecks & How to Fix Them {#bottlenecks}
| Profile Observation | Root Cause | Action |
|---|---|---|
| Attention memory-bound on roofline | HBM bandwidth ceiling at small batch | Enable FlashAttention; increase batch size |
| NCCL AllReduce gaps > 2ms/layer | Tensor-parallel communication overhead | Reduce TP degree; enable overlap |
| L2 hit rate < 20% on decode | KV cache exceeds L2 at long context | Enable prefix caching; larger block size |
| Prefill >> decode in timeline | No prefill-decode separation | Disaggregated inference routing |
| Python ops in critical CUDA path | Eager-mode dispatch overhead | torch.compile + CUDA Graphs |
| Low SM occupancy on GEMM | Small matrix dimensions (batch=1) | Continuous batching; group requests |
| Memory throughput < 40% despite memory-bound | PCIe bottleneck (CPU-GPU transfer in loop) | Move data to GPU before kernel loop |
| 100K tiny kernel launches | Kernel launch overhead dominates | Kernel fusion or CUDA Graphs |
| Warp divergence high | Branchy code within warps | Restructure to align thread execution |
| Register spilling detected | Too many registers per thread | Simplify kernel or use __launch_bounds__ |
11. Advanced: Tensor Core & Shared Memory Profiling {#advanced}
Tensor Core Investigation Checklist
1□ Are Tensor Core instructions executing? 2□ What is the achieved Tensor Core throughput? 3□ Is the workload compute-bound? 4□ Are dimensions suitable for WMMA/mma.sync? 5□ Is memory feeding the Tensor Cores fast enough?
Shared Memory Optimization Loop
1Profile → Bank conflicts detected 2 ↓ 3Change layout: tile[32][32] → tile[32][33] 4 ↓ 5Profile again → Conflicts reduced 6 ↓ 7Measure end-to-end speedup
Stall Reasons Deep Dive
Nsight Compute answers: "Why are warps waiting?"
Possible stall reasons:
1- Memory dependency 2- Execution dependency 3- Synchronization 4- Instruction latency 5- Not enough active warps 6- Pipeline limitations
12. FAQ {#faq}
Should I use nvprof in 2026?
No. nvprof and nvvp are deprecated since CUDA 11. Use nsys and ncu exclusively.
Can I profile on shared/cloud GPUs?
Most serverless platforms block hardware counter collection (ERR_NVGPUCTRPERM). Use bare-metal instances or VMs with SYS_ADMIN capability for ncu. nsys usually works everywhere.
How do I profile distributed training?
Use torch.profiler + HolisticTraceAnalysis for per-rank attribution and communication overlap. nsys can profile MPI jobs with per-rank output naming.
What's the difference between --set default and --set full?
default: One replay pass, basic roofline. Fast.full: Complete hardware counter collection. Slower but comprehensive.
How do I profile vLLM or TensorRT-LLM?
For vLLM: nsys profile --wait all python -m vllm.entrypoints.openai.api_server ...
For TensorRT-LLM: Build with --profiling-verbosity detailed and use nsys as the primary tool.
13. Conclusion {#conclusion}
Profiling transforms CUDA programming from art into engineering. By combining Nsight Systems for the big picture and Nsight Compute for kernel-level precision, you stop guessing and start optimizing with data.
Your action plan:
- Start with
nsys— Find the hot kernel and timeline bottlenecks - Drill with
ncu— Classify as memory-bound or compute-bound using the Roofline model - Apply the right fix — Coalescing for memory, Tensor Cores for compute, fusion for both
- Measure again — Close the Golden Profiling Loop
- Scale up — Use continuous batching, CUDA Graphs, and distributed profiling for production AI
"Don't optimize randomly. Profile first."
Related Articles:
- Multi-GPU Training with NCCL: From Single GPU to Distributed LLMs
- CUDA Kernel Optimization: Shared Memory, Occupancy & Tensor Cores
- FlashAttention Deep Dive: Memory-Efficient Transformer Inference
- torch.compile & CUDA Graphs: Eliminating Python Overhead
- Deploying LLMs with vLLM: A Production Engineering Guide