Nsight Compute Profiling

The CUDA profiler that tells you which pillar broke

Nsight Compute is NVIDIA's per-kernel profiler. It attaches to a running CUDA process, captures detailed counters for every kernel launch, and presents them in a readable dashboard. For BLAS work, the metrics that matter most are six:

Metric	Value (cuBLAS GEMM 4096³ on RTX 4090)	What it means
Compute (SM) Throughput	80.27%	Tensor pipes lit — that's the goal
Memory Throughput	42.54%	L2 is helping; DRAM half idle
DRAM Throughput	18.15%	GEMM is compute-bound, not BW-bound — confirmed
L2 Hit Rate	93.64%	Tile sizes keep operands hot in L2
Achieved Occupancy	16.64%	Each SM schedules ~8 warps (limited by regs + smem)
FMA pipeline busy	~73%	GEMM dominated by fused MUL-ADD instructions

Internal kernel duration: 2.63 ms ≈ 52 TF/s on the FP16-Tensor-Core path. Wall time ~2118 ms because Nsight injects 40+ instrumented profiling passes. Always read the 'internal' time, not the wall time, when judging perf inside Nsight.

Reading the picture

That table is what 'healthy' looks like. The kernel is compute-bound (80% SM), L2 is doing its job (94% hit rate), and the FMA pipe is the limiting factor. Headroom is essentially gone — this is a ~26%-from-perfect kernel that the library pillars have already maxed out.

What an unhealthy kernel looks like:

SM Throughput < 30% + DRAM Throughput > 80% → bandwidth-bound. You're not feeding the cores.
SM Throughput < 30% + DRAM < 30% → idle / launch-overhead dominated. Too few blocks, or kernel too short.
L2 Hit Rate < 50% → tile size doesn't fit. Reduce tile, or fuse the surrounding ops.
Occupancy < 10% → register spilling. Try --maxrregcount or restructure to use fewer registers.

Code

Run Nsight Compute on a CUDA binary·bash

# Capture the metrics shown in the table above:
ncu --set full ./gemm_bench --m 4096 --n 4096 --k 4096

# Lighter capture (just the BLAS-relevant metrics):
ncu --metrics sm__cycles_active.avg.pct_of_peak_sustained_elapsed,\
    dram__throughput.avg.pct_of_peak_sustained_elapsed,\
    lts__t_sector_hit_rate.pct,\
    sm__warps_active.avg.pct_of_peak_sustained_active \
    ./gemm_bench

# Profile only the kernel named 'gemm_tiled' (skip warmup launches):
ncu --kernel-name gemm_tiled ./gemm_bench

Programmatic Nsight start/stop ranges (when you have many launches)·cuda

#include <nvtx3/nvToolsExt.h>

// Mark a region you actually want to profile
nvtxRangePush("hot_loop");
for (int i = 0; i < 100; ++i) {
    cublasSgemm(...);     // your real workload
}
nvtxRangePop();

// Then capture only that range:
//   ncu --nvtx --nvtx-include 'hot_loop' ./your_app

Exercise

On a CUDA box, profile your tiled GEMM from Track 6 with ncu --set full. Capture the six metrics in the table above. Compare to cuBLAS's report on the same shape. The gap between your 'SM Throughput' and cuBLAS's 80% is exactly the room left to optimize. Then look at which pillar each gap reflects: if your DRAM Throughput is 80% and SM is 30%, you need better tiling. If your SM is 50% and DRAM is 20%, you need Tensor Core paths.

The CUDA profiler that tells you which pillar broke

Reading the picture

Code

External links

Exercise

Progress

Comments 0