Hands-on profiling workflow with Nsight Systems and Nsight Compute — finding bottlenecks, occupancy analysis, memory throughput.
This tutorial walks through NVIDIA's profiling tools end-to-end: from high-level timeline analysis with Nsight Systems, to deep kernel-level analysis with Nsight Compute, to interpreting metrics and fixing real bottlenecks.
Tutorials 01–05 (GPU Architecture, First Kernel, Memory & Bandwidth, Thread Scheduling, Shared Memory & Tiling). You should be comfortable writing and launching CUDA kernels before learning to profile them.
The gap between "premature optimisation" and "informed optimisation" is measurement. Profiling tells you exactly where your kernel spends its time — and whether you're limited by compute, memory, or something else entirely.
The roofline model plots achievable performance (GFLOP/s) against operational intensity (FLOP/byte). It reveals whether a kernel is compute-bound or memory-bound — the single most important classification for optimisation strategy.
If your kernel sits on the slope, you're memory-bound — optimise memory access patterns (coalescing, caching, tiling). If it sits on the flat ceiling, you're compute-bound — reduce arithmetic, use intrinsics, or leverage Tensor Cores.
Nsight Systems is your starting point for profiling. It captures a system-wide timeline showing CPU activity, CUDA API calls, kernel launches, and memory transfers — giving you the big picture before you dive into individual kernels.
Nsight Systems ships with the CUDA Toolkit. Verify your installation:
# Check nsys is installed
nsys --version
# If not found, install via CUDA Toolkit or standalone package
sudo apt install nsight-systems # Ubuntu/Debian# Basic profiling — generates a .nsys-rep report file
nsys profile --stats=true -o my_report ./my_app
# With trace options for CUDA and OS runtime
nsys profile \
--trace=cuda,osrt,nvtx \
--cuda-memory-usage=true \
--stats=true \
-o matmul_profile \
./matmul
# Export to SQLite for custom analysis
nsys export --type=sqlite my_report.nsys-repThe Nsight Systems GUI shows parallel rows for each activity stream. Here is what a typical timeline looks like:
Empty space = idle GPU. Usually caused by synchronisation, CPU bottlenecks, or insufficient overlap.
Large transfers between H→D or D→H relative to kernel time suggest you should use pinned memory or streams.
Kernel launch overhead (~5–10 μs) dominates if the kernel itself runs for only microseconds. Consider kernel fusion.
Always start with nsys profile for the big picture. Only move to Nsight Compute once you know which kernel to optimise.
Nsight Compute (ncu) is the kernel-level profiler. It replays your kernel multiple times to collect detailed hardware performance counters, giving you a precise breakdown of what the hardware is doing during execution.
# Full metric set for all kernels
ncu --set full ./my_app
# Profile a specific kernel by name
ncu --set full \
--kernel-name matmul_kernel \
--launch-skip 0 \
--launch-count 1 \
./my_app
# Save report for GUI analysis
ncu --set full -o matmul_report ./my_app
# Open in GUI
ncu-ui matmul_report.ncu-rep
# Compare two reports (before/after optimisation)
ncu --set full -o before ./my_app_v1
ncu --set full -o after ./my_app_v2
ncu-ui --diff before.ncu-rep after.ncu-repShows what % of theoretical peak you achieved for both compute and memory. The higher number tells you the bottleneck.
Executed operations vs peak. Broken down by pipe: FMA, ALU, FP64, Tensor, etc. Low values mean wasted cycles.
Bytes moved vs peak bandwidth for each memory level: global, L2, L1/shared. Reveals cache effectiveness.
Theoretical max warps vs achieved. Shows limiting factors: registers per thread, shared memory per block, or block size.
The SOL chart is the first thing to check. It shows compute and memory utilisation as a percentage of the hardware peak:
Nsight Compute replays each kernel multiple times to collect counters. This means profiling is much slower than normal execution. Use --launch-count and --kernel-name to limit which kernels are profiled.
Occupancy is the ratio of active warps on an SM to the maximum number the SM supports. Higher occupancy generally helps hide memory latency — but 100% is not always the goal. The right level depends on your kernel's characteristics.
The maximum possible given your kernel's resource usage (registers, shared memory, block size). Computed statically before the kernel runs.
The actual average number of active warps divided by the maximum, measured during execution. Always ≤ theoretical occupancy.
Three resources compete for SM capacity. Whichever one runs out first limits your occupancy:
| Limiting Factor | How It Limits Occupancy | How to Check | Typical Fix |
|---|---|---|---|
| Registers per thread | SM has 65,536 registers. If each thread uses 64, a block of 256 threads needs 16,384 — only 4 blocks fit. | ncu --set full → Occupancy section |
__launch_bounds__, reduce local variables, use -maxrregcount |
| Shared memory per block | SM has 48–228 KB shared memory. Large tile sizes can consume it all, limiting concurrent blocks. | ncu --set full → Occupancy section |
Reduce tile size, use dynamic shared memory, tune with cudaFuncSetAttribute |
| Block size (threads per block) | SM supports max 1,536–2,048 threads. Large blocks (1024) mean fewer concurrent blocks per SM. | ncu --set full → Occupancy section |
Experiment with 128, 256, or 512 threads per block |
// Query max active blocks for a given kernel and block size
int numBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&numBlocks,
matmul_kernel, // kernel function
256, // block size
0 // dynamic shared memory bytes
);
printf("Max active blocks per SM: %d\n", numBlocks);
// Let CUDA suggest an optimal block size
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(
&minGridSize,
&blockSize,
matmul_kernel, // kernel function
0, // dynamic shared memory per block
0 // block size limit (0 = no limit)
);
printf("Suggested block size: %d\n", blockSize);
printf("Min grid size for full occupancy: %d\n", minGridSize);High occupancy does not guarantee high performance. A kernel at 50% occupancy with efficient memory access can outperform one at 100% occupancy with poor coalescing. Use occupancy as one input, not the only metric.
Most CUDA kernels are memory-bound. Understanding where your bytes travel — and how efficiently — is critical for getting close to the roofline. Nsight Compute breaks down memory throughput at every level of the hierarchy.
Global memory efficiency measures how many of the bytes fetched from global memory were actually requested by the kernel. Uncoalesced accesses waste bandwidth by loading unused bytes.
32 threads in a warp each access consecutive 4-byte elements. One 128-byte transaction serves all 32 threads.
32 threads each access every 4th element. Four 128-byte transactions needed, but only 128 bytes are useful.
Cache hit rates tell you how well your access pattern exploits spatial and temporal locality:
# Query specific memory metrics
ncu --metrics \
l1tex__t_sectors_pipe_lsu_mem_global_op_ld_hit_rate.pct,\
lts__t_sectors_srcunit_tex_op_read_hit_rate.pct,\
smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct \
./my_app| Metric | What It Measures | Good Value | Poor Value Indicates |
|---|---|---|---|
| Global Load Efficiency | Useful bytes / total bytes fetched for loads | > 80% | Uncoalesced / strided access |
| Global Store Efficiency | Useful bytes / total bytes for stores | > 80% | Uncoalesced / scattered writes |
| L1 Hit Rate | Fraction of L1 requests served without going to L2 | > 50% | Poor spatial locality, large working set |
| L2 Hit Rate | Fraction of L2 requests served without going to DRAM | > 60% | Working set exceeds L2 capacity |
| DRAM Throughput | Bytes/sec actually moved from/to global memory | > 70% of peak | Access pattern issues, low occupancy |
# Sectors per request > 1 means uncoalesced access
ncu --metrics \
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio,\
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_st.ratio \
./my_app
# Ideal: ratio close to 1.0 for 32-bit loads
# Values of 4+ indicate severe coalescing problemsIf global load efficiency is below 50%, fixing your memory access pattern (coalescing, shared memory tiling, or data layout changes) will almost certainly give a bigger speedup than any compute optimisation.
Let's walk through profiling two versions of matrix multiplication: a naive implementation and a shared-memory tiled version. We'll see exactly how the metrics change and why.
# First, get the big picture with Nsight Systems
nsys profile --stats=true -o naive_timeline ./matmul_naive
# Then drill into the kernel with Nsight Compute
ncu --set full \
--kernel-name matmul_naive_kernel \
--launch-count 1 \
-o naive_kernel \
./matmul_naivensys profile --stats=true -o tiled_timeline ./matmul_tiled
ncu --set full \
--kernel-name matmul_tiled_kernel \
--launch-count 1 \
-o tiled_kernel \
./matmul_tiledHere's what the metrics look like side by side for a 1024×1024 single-precision matrix multiply on an RTX 4090:
| Metric | Naive | Tiled (32×32) | Change |
|---|---|---|---|
| Kernel Duration | 4.82 ms | 0.61 ms | 7.9× faster |
| SOL SM [%] | 12.1% | 68.4% | ↑ 5.7× |
| SOL Memory [%] | 78.6% | 34.2% | ↓ (good: no longer the bottleneck) |
| DRAM Throughput | 742 GB/s | 108 GB/s | 6.9× less DRAM traffic |
| L1 Hit Rate | 0% | 0% | — (shared mem bypasses L1) |
| Achieved Occupancy | 49.8% | 49.2% | — (similar) |
| Global Load Efficiency | 100% | 100% | — (both coalesced) |
| Shared Memory Throughput | 0 GB/s | 3,412 GB/s | Tiling effect |
Bottleneck: DRAM bandwidth
Bottleneck: shifting toward compute
Tiling reduced DRAM traffic by ~7× by converting redundant global loads into fast shared memory reads. The kernel moved from memory-bound to compute-bound on the roofline — now further gains require compute optimisations (loop unrolling, vectorised loads, Tensor Cores).
After profiling many CUDA kernels, the same patterns emerge repeatedly. This reference table maps observable symptoms in profiler output to their likely root causes and fixes.
| Symptom in Profiler | Likely Cause | Fix |
|---|---|---|
| SOL Memory >> SOL Compute | Kernel is memory-bound; spending most time waiting for data | Shared memory tiling, data reuse, coalescing, cache-friendly layout |
| SOL Compute >> SOL Memory | Kernel is compute-bound; ALUs are saturated | Algorithmic improvements, Tensor Cores, reduced precision (FP16/BF16) |
| Both SOL Compute and SOL Memory low | Latency-bound: not enough active warps to hide memory latency | Increase occupancy (tune block size, reduce register pressure) |
| Global load efficiency < 50% | Uncoalesced memory access (strided or random pattern) | Restructure data layout (AoS → SoA), transpose access, use shared memory |
| High DRAM traffic but low L2 hit rate | Working set too large for L2 cache | Tile the algorithm, process data in cache-sized chunks |
| Achieved occupancy << theoretical | Block-level synchronisation stalls, load imbalance, tail effect | Balance workload per block, reduce __syncthreads() frequency |
| Long gaps between kernels in timeline | CPU bottleneck or synchronisation between launches | Use CUDA streams, overlap compute and memory, reduce host-side work |
| H→D / D→H transfers dominate timeline | Excessive data transfer overhead | Pinned memory (cudaMallocHost), async transfers, keep data on GPU |
| High warp stall (not selected / memory dependency) | Memory latency not hidden by available parallelism | Increase ILP (instruction-level parallelism), prefetch, increase occupancy |
| High branch divergence in warp stall reasons | Threads within a warp taking different paths | Reorganise data so nearby threads take same path, predication |
Never optimise without profiling first. Never assume your fix worked without profiling again. The profiler is the source of truth.
nsys profilencu --set fullNsight Systems gives the big picture. Nsight Compute gives the detail. Always start with the former, then drill into the hottest kernel with the latter.
The SOL chart is your compass. Memory-bound and compute-bound kernels need completely different optimisation strategies. Applying the wrong one wastes time.
# System-wide timeline
nsys profile --stats=true -o report ./my_app
# Kernel deep dive
ncu --set full -o kernel_report ./my_app
# Compare before/after
ncu-ui --diff before.ncu-rep after.ncu-rep
# Query specific metrics
ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
dram__throughput.avg.pct_of_peak_sustained_elapsed \
./my_appCUDA Streams & Concurrency — overlap kernel execution with memory transfers, launch multiple kernels concurrently, and use events for fine-grained timing and synchronisation.