4 of 4

Profiling, benchmarking, and correctness

A blocked or tiled program is, by design, no longer readable as a literal description of what one thread does. That means neither speed nor correctness can be checked by re-reading the source the way the per-thread CUDA code of Foundations could be: both need to be measured against what actually ran on the GPU. NVIDIA splits those two jobs across separate tools. Nsight Compute answers “how fast, and where is the time going”; Compute Sanitizer answers “did it read or write anything it shouldn’t have.”

Nsight Compute profiles by inserting measurement libraries into the running application process, intercepting its communication with the CUDA driver, and collecting metrics whenever it detects a kernel launch. Its most direct performance readout is occupancy: “occupancy is the ratio of the number of active warps per multiprocessor to the maximum number of possible active warps,” and a large gap between the theoretical and the achieved value “typically indicates highly imbalanced workloads.” Its GPU Speed Of Light section reports the achieved percentage of compute and memory throughput against the theoretical maximum for each, which is the same flop-bound-versus-memory-bound question from Foundations, now measured per kernel instead of estimated from a datasheet.

Collecting all of that is not free, and not always possible in one pass: the GPU has a limited number of hardware counters it can read concurrently, and some metrics require patch-based software counters whose overhead would itself distort the measurement. When a requested set of metrics cannot be collected together, Nsight Compute uses kernel replay , saving all memory the kernel can reach before the first pass and restoring whatever it wrote before each subsequent one. A profiled run is therefore not the same execution as an unprofiled one, which is a reason to treat wall-clock numbers taken while profiling as approximate.

One command, four checkers. Compute Sanitizer runs the same kernel under one of four instrumented modes, each watching for a failure class a profiler never checks: bad addresses, shared-memory races, uninitialized reads, and misused barriers.Illustrative numbers

Compute Sanitizer is the correctness half, shipped as part of the CUDA toolkit. Its own documentation states the reason it exists plainly: “every programmer invariably encounters memory access errors and thread ordering hazards that are hard to detect and time consuming to debug,” and “the number of such errors increases substantially when dealing with thousands of threads.” None of its four tools measure speed at all; they instrument actual memory and synchronization behavior and report a violation the moment one occurs, which is precisely the check a profiler has no reason to perform. A kernel can report high occupancy and near-peak memory throughput in Nsight Compute while still failing racecheck, because occupancy says nothing about whether two warps are racing on the same shared-memory address.

Sections, and what each one answers#

Nsight Compute does not collect everything by default. It “uses Section Sets (short sets) to decide, on a very high level, the number of metrics to be collected. Each set includes one or more Sections, with each section specifying several logically associated metrics.” The basic set runs when no --set, --section, and no --metrics option is passed, and it is deliberately cheap: mostly “high-level utilization information as well as static launch and occupancy data,” the latter two “regularly available without replaying the kernel launch.” Everything beyond it costs replay passes, which is why --set full is a decision rather than a default.

The section catalogue reads as a list of performance questions. GPU Speed Of Light Throughput is the one to read first: a “high-level overview of the throughput for compute and memory resources of the GPU. For each unit, the throughput reports the achieved percentage of utilization with respect to the theoretical maximum. Breakdowns show the throughput for each individual sub-metric of Compute and Memory to clearly identify the highest contributor.” That is the flop-bound versus memory-bound question of Foundations asked of one kernel instead of one datasheet, and the breakdown is the actionable half, because it names the specific sub-metric sitting closest to its own ceiling.

Memory Workload Analysis then splits the memory answer into three distinct bottlenecks. Memory “can become a limiting factor for the overall kernel performance when fully utilizing the involved hardware units (Mem Busy), exhausting the available communication bandwidth between those units (Max Bandwidth), or by reaching the maximum throughput of issuing memory instructions (Mem Pipes Busy).” Saturated units, saturated links, and a saturated issue rate call for different fixes, and the section’s chart is built to tell them apart: links from the kernel to the logical memory spaces are labelled with executed instructions, links into the physical units with the requests those instructions generate, and every link is colored by percentage of peak. Ports are colored separately, for the case the link numbers alone would hide: “While the links sharing a port might operate well below their individual peak performances, the unit’s data port may have already reached its peak.” The tables underneath add per-unit detail, including a bank-conflict count for shared memory, and one warning worth internalizing: high utilization “can show potential bottlenecks, as it does not necessarily indicate efficient usage.” A cache saturated by uncoalesced requests reads as busy.

The roofline section closes the loop with Foundations directly. It plots the kernel as a single achieved point against two rooflines, with floating-point operations per second on the vertical axis and “Arithmetic Intensity, which is the ratio between Work (expressed in floating point operations per second), and Memory Traffic (expressed in bytes per second)” on the horizontal, and the ridge point where the sloped memory boundary meets the flat compute one. “The distance from the achieved value to the respective roofline boundary” is the available headroom, and the reading rule that decides what to do next is stated outright: “An achieved value that lies on the Memory Bandwidth Boundary but is not yet at the height of the ridge point would indicate that any further improvements in overall FLOP/s are only possible if the Arithmetic Intensity is increased at the same time.” A kernel already sitting on the sloped boundary will not be rescued by better loads. It has to move right, which means fusing work or reusing more data per byte fetched, and that is a change to the algorithm rather than to the schedule.

Beyond a single kernel#

Nsight Compute answers its questions one kernel at a time, and that is also its blind spot: it cannot say whether the kernel it is dissecting is the one worth dissecting. That is Nsight Systems’s job. By default it “collects a profile over the entire run of your application,” laying out CPU threads, CUDA API calls, and GPU activity on one timeline, and its guide recommends narrowing collection to the performance-critical region with cudaProfilerStart() and cudaProfilerStop() rather than profiling a test harness’s setup and validation. NVTX markers and ranges added to the application appear in the Timeline View and are projected onto the GPU timeline, “allowing you to see what GPU activity was launched within each CPU range.” The workflow this implies runs in one direction: find the expensive kernel, or the gap where no kernel is running, on the Systems timeline first, then drill into that kernel with Compute.

Even with the right kernel in hand, a number measured casually is not a number worth reporting. NVIDIA’s GEMM performance measurement methodology, published with the CUTLASS documentation, prescribes what a reproducible benchmark harness must do: separate warmup and profiling loops, delimited by cudaProfilerStart/Stop or CUDA events; no allocations, copies, or extra kernels between launches; and buffer rotation, cycling each tensor through duplicate buffers whose total footprint is at least twice the L2 capacity, so every iteration starts from DRAM instead of inheriting the last iteration’s cache. It distinguishes fixed frequency tests, which measure architectural and software efficiency at a locked clock, from fixed power tests, which mimic the dynamically-scaled clocks of real-world use and vary far more. The scale it recommends is sobering for anyone timing a kernel in a loop of ten: earlier GEMM studies needed more than 1,000 iterations for stable results on a 4096 by 4096 by 4096 problem, and for large GEMMs on Blackwell the document uses 10,000 warmups, 4,000 profiling iterations, and a second of cool-down between tests.

A harness built to that spec is still only a claim until something checks it, and the document supplies the check. Before any data analysis, three profiled metrics validate that a benchmark followed its own rules. The first closes the loop on buffer rotation: use “the L2 hit rate per-iteration to ensure that buffers are rotated appropriately,” and with Nsight Compute do it under --cache-control none so that the profiler’s own cache flushes, described below, are not what is making the cache look cold. Rotation you cannot see in the hit rate did not happen. The second is the profiled GPC frequency, examined for stability across the profiling iterations, which is not a formality: running GEMMs back to back, the clocks “were observed to oscillate for around the first 3 seconds of runtime before settling on a stable frequency.” The third is the kernel launch times, examined for minimal gaps between iterations.

Only then is there a number, and the document wants it compared rather than reported. The measured average runtime goes against the speed-of-light runtime at the locked frequency for a fixed frequency test, or at the settled average clock frequency for a power-constrained one, which is the roofline of Foundations returning as an acceptance criterion instead of a diagram. Even the fill pattern follows from this: fixed frequency tests “are not sensitive to the data fill pattern,” so the recommendation is zero-fill, which draws less power and makes the locked clock easier to sustain, which is what keeps the frequency check in the paragraph above from failing.

AMD’s counterpart to the kernel drill-down is ROCm Compute Profiler, also known by its package name rocprofiler-compute: “a kernel-level profiling tool for machine learning and high performance computing (HPC) workloads running on AMD Instinct GPUs.” It is built on ROCprofiler-SDK to monitor hardware performance counters, acquires those counters via application replay, and runs accelerator-specific microbenchmarks to build hierarchical roofline data. Its analysis vocabulary maps almost term for term onto the NVIDIA one: system Speed-of-Light and hardware-block-level Speed-of-Light summaries, memory chart and roofline analysis, and baseline comparisons between runs, all for the CDNA GPUs whose waves and chiplets HipKittens schedules around in the previous section.

The command lines#

All of that reduces to a small number of invocations. Nsight Systems is the wide pass and Nsight Compute the narrow one, and both documentations give their example command lines directly. The following are taken from those pages, with the angle-bracket placeholders they use left in:

Collection commands, from the Nsight Systems and Nsight Compute documentation

nsys profile --trace=cuda,nvtx -d 20
    --sample=none --cpuctxsw=none -o my_test <application>
    [application-arguments]

nsys stats --report cuda_gpu_trace --report cuda_gpu_kern_sum --report cuda_api_sum --format csv,column --output .,- report1.nsys-rep

ncu --set full --metric-distribution-groups 4 --communicator none -o report <app> [app arguments]

ncu --target-processes all -o <report-name> mpirun [mpi arguments] <app> [app arguments]

ncu --nvtx --nvtx-include "A_range/" CuNvtx.exe

The first line is a deliberately narrow Systems collection. --trace=cuda,nvtx limits tracing to CUDA and NVTX, -d 20 ends collection after 20 seconds or at application exit, whichever comes first, and --sample=none --cpuctxsw=none switches off CPU instruction-pointer sampling and thread scheduling. Plain nsys profile <application> instead traces “CUDA, OpenGL, NVTX, and OS runtime libraries APIs” and collects both of the things the narrow form turned off, which is more than a kernel investigation needs. When the interesting region is a phase rather than the whole run, --capture-range accepts cudaProfilerApi, nvtx, or a hotkey, so “profiling will start only when an appropriate start API or hotkey is invoked,” and --capture-range-end decides whether the session then stops, shuts down, or repeats for a bounded number of ranges.

The second line turns a report into tables without opening the user interface. nsys stats exports an SQLite database beside the report and runs named report scripts over it: cuda_gpu_kern_sum and cuda_api_sum answer which kernel and which API call are eating the run, cuda_gpu_trace is the per-launch list underneath them, and the paired --format and --output lists send the first report to a CSV file and print the other two as columns. Passing --stats=true to nsys profile generates a default fixed set of summary statistics inline at the end of a run.

The Compute lines are the opposite posture. --set full asks for every section, and the two flags beside it are the Metric Distributor: the metrics are divided into four groups collected by four processes across the machine’s GPUs, which reduces the number of passes and leaves four partial reports to merge. Narrowing is the usual next move: -k filters kernels by name or by a regex: expression, -c limits how many launches are collected and -s skips launches before collection starts, --kernel-id filters by context, stream, name, and invocation number, and --section names individual sections in place of a whole set. Output is controlled separately, through --page for details, raw, source, or session, --csv to make any of them machine readable, and --print-source to choose between SASS, PTX, CUDA, and a CUDA-to-SASS correlation. --target-processes all is what makes the mpirun form work, and because one report per rank is needed, the guide writes the output name through a macro such as %q{OMPI_COMM_WORLD_RANK}. The last line applies the same filtering through NVTX rather than kernel names.

What none of these lines say out loud is how much they change the run. Nsight Compute serializes kernel launches within the profiled application and across processes, because only one process can profile a given device at a time. It “attempts to limit GPU clock frequencies to their base value” so a metric does not depend on where in the application the kernel happened to sit, adjustable with --clock-control. And it flushes all GPU caches before each replay pass so every pass sees a clean cache, which --cache-control none disables when the measurement is deliberately about a warm one. Add kernel replay’s save and restore of device memory, and a profiled execution is not the execution being shipped, which is why the benchmark methodology above insists on a separate harness for any number reported as a speed.

Correctness has a command line too, and it is shorter. Compute Sanitizer wraps the application the same way the profilers do, and a --tool value selects which of the four checkers in the figure above actually runs: memcheck, the default, for out-of-bounds and misaligned accesses to global, local, and shared memory, plus hardware-reported errors and memory leaks; racecheck for shared-memory data-access hazards that can cause data races; initcheck for uninitialized accesses to device global memory; and synccheck for invalid usage of synchronization primitives. One tool runs per invocation, so a clean bill of health is four runs, not one.

Compute Sanitizer runs and the compilation flags they want

compute-sanitizer [options] app_name [app_options]

compute-sanitizer --tool memcheck --leak-check full <app> [app arguments]
compute-sanitizer --tool racecheck --racecheck-report all <app> [app arguments]
compute-sanitizer --tool initcheck <app> [app arguments]
compute-sanitizer --tool synccheck <app> [app arguments]

nvcc -Xcompiler -rdynamic -lineinfo -o out in.cu

nvcc -fdevice-sanitize=memcheck -lineinfo -o out in.cu

The last two lines are the half that gets skipped. The tools “do not need any special compilation flags to function,” which is exactly why a first run is so often useless: without line information a hazard is reported against an address rather than against a line of the kernel that caused it. -lineinfo generates line-number information “without affecting the optimization level of the output,” -G generates full debug information at the cost of that optimization, and the tools “can display source attribution of errors for applications compiled with line information.” On Linux the host compiler wants -rdynamic, passed through -Xcompiler, so that host backtraces keep their function names. Memcheck has a second, faster path, -fdevice-sanitize=memcheck, which instruments at compile time instead of at run time and adds base-and-bounds analysis, and the documentation flags the trap in it: that option “does not imply the generation of debug information,” so -lineinfo or -G is still needed alongside it.

The per-tool options narrow a run the way -k narrows an ncu one. --leak-check full prints every allocation never freed through cudaFree by the time the context was destroyed. --racecheck-report chooses between hazard, analysis, and all, with analysis the default. --initcheck-address-space widens initcheck from global memory to shared or both. And --check-warpgroup-mma, on sm_90a, extends memcheck and synccheck to PTX wgmma instructions, with memcheck checking that the matrices loaded by wgmma.mma_async are in shared-memory range: the same instruction family Kernel optimization measured a Hopper matmul against, now being checked rather than timed.