Profiling with Nsight: reductions and transposes

1. Two tools, two questions

Everyone who starts optimizing GPU code goes through the same phase: change something, rerun the benchmark, stare at a number that moved by 3% in the wrong direction, and have no idea why. The benchmark tells you that something is slow. It can never tell you why. For that you need a profiler, and NVIDIA ships two, because there are two different questions:

The order matters: always start with the timeline. Optimizing a kernel that takes 5% of the run, while the GPU idles 40% of the time between kernels, is the most common way to waste a week in this field. Nsight Systems tells you whether any kernel is worth opening; Nsight Compute tells you what is wrong with it.

2. Profiling under WSL2

Both tools are part of the CUDA toolkit you installed in chapter 1, and both work inside WSL2. NVIDIA supports profiling there on Windows 11 with driver 525 or newer (Windows 10 needs build 19044 and driver 545). The pieces fit together like this:

• The command-line tools run inside WSL. nsys profile and ncu run your program and write a report file (.nsys-rep, .ncu-rep). Both also print text summaries, which is often all you need.
• The GUIs run on Windows. Install Nsight Systems and Nsight Compute for Windows and open the reports from \\wsl$\Ubuntu-24.04\home\…. Reports are portable: a report made anywhere opens anywhere.
• Counter permission is a Windows setting. The Windows driver owns the GPU’s performance counters, so the switch from chapter 1 (NVIDIA Control Panel → Developer → Manage GPU Performance Counters → all users) is what makes ncu work. Without it you get ERR_NVGPUCTRPERM.

Three WSL-specific caveats, all from NVIDIA’s own documentation:

• Nsight Compute’s context-switch trace (part of its timeline-style PM sampling) is not supported on WSL. Nothing in this course needs it.
• Nsight Systems’ 2025 release notes list a known issue on WSL: its default GPU timestamp conversion is unreliable, so GPU work can appear shifted relative to the CPU calls that caused it. The documented workaround is one line in its config file, and it is in lab 4’s README. Apply it only if your timeline looks wrong.
• We trace only CUDA and NVTX, and switch off Nsight Systems’ CPU sampling and OS context-switch tracing. Those rely on Linux kernel performance features that a WSL2 kernel may not fully expose, and nothing in this chapter needs them.

3. Nsight Systems: reading a timeline

A timeline has one row per thing that can be busy. The ones you will read every day:

• CUDA API (per CPU thread): every call into the runtime, as a bar as long as the call took on the CPU. cudaLaunchKernel bars are short; cudaDeviceSynchronize and cudaMemcpy bars are as long as the CPU was blocked.
• GPU kernels (per stream): every kernel as a bar as long as it ran on the GPU, with a line back to the launch call that created it.
• Memory copies: host-to-device and device-to-host transfers, with their size and speed.
• NVTX: ranges you put in your own code with nvtxRangePushA("name") / nvtxRangePop(). They are how you find “the part where we do X” in a 30-second trace of a real engine. PyTorch, vLLM and llama.cpp all emit NVTX ranges.

What you look for is almost always the same thing: white space in the GPU row. Every gap is time the GPU could have worked and did not, and the rows above tell you why. Lab 4’s graphs.cu builds the three most common situations on purpose: 1,000 tiny kernels (4,096 floats each, a couple of microseconds of GPU work) launched three different ways.

CUDA graphs

A CUDA graph is a recording of GPU work that can be replayed with one call. You make one by capturing a stream: between cudaStreamBeginCapture and cudaStreamEndCapture, launches are recorded instead of executed. cudaGraphInstantiate then validates and prepares the whole sequence once (the expensive step), and cudaGraphLaunch replays it as often as you like:

cudaStreamBeginCapture(st, cudaStreamCaptureModeGlobal);
for (int k = 0; k < 1000; ++k) bump<<<16, 256, 0, st>>>(d, 4096);   // recorded, not run
cudaStreamEndCapture(st, &graph);
cudaGraphInstantiate(&exec, graph, 0);   // once
cudaGraphLaunch(exec, st);                // 1,000 kernels, one call

The catch is that a graph replays exactly what was recorded: same kernels, same grid sizes, same pointers. That fits workloads that repeat the same shape over and over, which is precisely LLM decode: every step runs the same few hundred kernels on the same buffers. This is why inference engines capture graphs for decode (vLLM captures one per batch size it expects; llama.cpp uses them on NVIDIA GPUs), and it is how your spine project will fix the overhead problem that chapter 1 predicted for GPT-2 small.

From picture to numbers

The GUI is for looking; nsys stats is for numbers. make nsys in lab 4 prints three tables from the report: nvtx_sum (time per labelled range), cuda_api_sum (how many times each API call happened and its average CPU time: 2,000 cudaStreamSynchronize calls in A is a smell you can see without opening a timeline), and cuda_gpu_kern_sum (total GPU time per kernel name, the first place to look in a real program).

One flag deserves a note: by default Nsight Systems draws a whole graph as a single bar. The lab passes --cuda-graph-trace=node so that you can see the individual kernels inside it.

4. Nsight Compute: one kernel under the microscope

How it works, and why its times differ from yours

A GPU can only count a limited number of events at once, but a full report needs thousands of metrics. So Nsight Compute replays each profiled kernel: it saves the memory the kernel touches, runs it, restores the memory, runs it again with a different set of counters, and repeats, dozens of times for the full set. Three consequences:

• Profiling is slow. Profile a few kernel launches, not a whole benchmark. That is why lab 4’s programs have a profile mode that runs each kernel exactly once.
• Caches are flushed between passes by default (--cache-control all), so every pass starts cold. A kernel that benefits from data left in L2 by the previous kernel looks slower under ncu than in your benchmark.
• Clocks are locked. Since Nsight Compute 2026.1 the default (--clock-control boost) locks the GPU to its boost clock; earlier versions locked it to the base clock. Your benchmarks run with whatever clock the card chooses at that moment. --clock-control none leaves the clocks alone.

So use ncu’s duration to compare kernels with each other inside one report, and your lab benchmark (or Nsight Systems, which neither locks clocks nor flushes caches) for real-world time.

Sets and sections

Results are organized in sections, and sections are grouped in sets. Without flags you get the basic set (Speed of Light, Launch Statistics, Occupancy, Workload Distribution): cheap, and enough to classify a kernel. --set full collects everything, including the sections that explain why:

Every section ends with rules: short diagnoses such as “uncoalesced shared accesses” or “this kernel exhibits low compute throughput and memory bandwidth utilization,” each with an estimated speedup. They are good hints and not always right: treat them as hypotheses to check, not verdicts.

Speed of Light: the one panel to read first

Metric names, decoded

Every number in a report is a named metric, and the names follow a grammar. Once you can read them you can request exactly what you want with --metrics, as lab 4’s make ncu-table does:

dram__throughput.avg.pct_of_peak_sustained_elapsed
└┬─┘  └───┬────┘ └┬┘ └──────────────┬─────────────┘
 unit   counter  rollup           suffix
(DRAM)  (what)   (over all DRAM   (as % of the peak rate,
                  channels)        over the kernel's run)

Units you will meet: gpu (whole chip), sm and smsp (an SM, and one of its four sub-partitions: the quadrants of chapter 2), l1tex (L1 and shared memory), lts (L2 slices), dram. The rollup says how per-unit values combine (.sum, .avg, .max), and ratios end in .ratio. The ones worth knowing by heart:

5. Stall reasons: what the warps were waiting for

Chapter 2 described the scheduler: every cycle, each quadrant picks one ready warp and issues its next instruction. Nsight Compute watches this happen. For every warp that could not issue, it records why, and reports the result as an average number of warps stalled on each reason per issued instruction. That table is the most direct answer a profiler can give to “why is this slow?”: it names the thing everyone is waiting for.

Here is that bookkeeping as a machine: the SM’s four quadrants of warps, a sampler that taps one warp at a fixed interval of cycles, and one counter per state. Run it on each kernel and watch which counter pulls ahead.

Two rules make stall reasons useful instead of confusing:

1. Read them together with Speed of Light. A kernel at 90% of DRAM bandwidth should be dominated by long_scoreboard (waiting for memory): that is what being memory-bound looks like, and there is nothing to fix. The same stall with DRAM at 30% means the kernel is waiting on latency it fails to hide, and you have work to do.
2. Stalls are charged to the instruction that waits, not the one that caused the wait. A slow load shows up on the first instruction that uses the loaded value. On the Source page, look a few lines up from the hot line.

6. The reduction, six versions

A reduction turns an array into one value: a sum, a maximum, a norm. It is everywhere in ML (softmax needs a max and a sum per row, layernorm a mean and a variance, every loss is a sum), and it is the standard exercise for learning to read a profiler, because the naive GPU version is slow for reasons that are all visible in counters.

Sequentially it is one loop. In parallel it is a tree: add pairs, then pairs of pairs, and so on. N values take N−1 additions, the same as the loop, but in log₂N steps instead of N. Rearranging the additions this way is only allowed because addition is associative and commutative. For integers that is exactly true. For floats it is almost true: each ordering rounds differently, so a GPU sum and a CPU sum of the same floats differ in the last bits, and two runs of a kernel that finishes with atomics can differ from each other. Lab 4’s test checks floats against a double-precision reference with a tolerance for this reason, and prints one surprise: a plain sequential float loop is the least accurate of all at large sizes, because once the running sum is big, every small addend loses most of its bits. Trees add numbers of similar size.

All six versions below reduce each block’s share in shared memory and finish with one atomicAdd per block into the final result (blocks cannot wait for each other, chapter 2). Each version changes one thing, and for each I state what the profiler should show. These are predictions from the mechanisms and from the compiled SASS, not measurements: make ncu-table in lab 4 prints the real values for all six on your card.

The first three are the textbook sequence (PMPP chapter 10, and Mark Harris’s famous 2007 slides). They were designed to teach divergence, and on 2007 hardware each step was worth a factor of two. The profiler is about to tell a more interesting story, so look at what the compiler does with v1 first. Because the block size is a compile-time constant, the loop is fully unrolled, every % becomes a bit mask, and the if becomes predication, not a branch (chapter 2): every warp issues every level’s instructions, with some lanes switched off:

LOP3.LUT P0, RZ, R2, 0x1, RZ, 0xc0, !PT   // P0 = (tid & 1) != 0: the "% 2", as a mask
BAR.SYNC.DEFER_BLOCKING 0x0               // __syncthreads()
@!P0 LDS R4, [R3+0x4]                     // s[tid + 1], only lanes with P0 false
@!P0 LDS R5, [R3]                         // s[tid]
@!P0 STS [R3], R4                         // (FADD omitted)
LOP3.LUT P0, RZ, R2, 0x7, RZ, 0xc0, !PT   // next level's mask, 0x3, 0x7, 0xf ...

v2 and v3 compile to the same shape: 88 instructions each, identical in count. So the three versions issue the same number of warp instructions, and the classic “v1 wastes issue slots on divergent warps” argument does not apply to this code. What still differs is which lanes are switched on, and that decides how much shared-memory traffic each level causes. A warp whose lanes are all off sends no request; a warp with even one active lane sends one, and a request with bank conflicts takes several passes (“wavefronts”). Step through it:

v1: interleaved addressing

for (unsigned stride = 1; stride < 256; stride *= 2) {
    if (tid % (2 * stride) == 0) s[tid] += s[tid + stride];
    __syncthreads();
}

The direct translation of the tree. The adding threads are the multiples of 2×stride, spread across the whole block, so for the first five levels every one of the eight warps has at least one active lane and sends its own nearly empty shared-memory request. Per block: 94 shared-load requests over the whole tree, each one a single pass (no conflicts: the active lanes of a warp sit on different banks). Expected evidence: smsp__inst_executed_op_shared_ld_pred_on_any.sum (shared loads with at least one lane on) about 4× higher than v3’s, and zero bank conflicts. Stall reasons: barrier dominates, as in every version with eight barriers per 256 values.

v2: compact thread index

unsigned idx = 2 * stride * tid;
if (idx < 256) s[idx] += s[idx + stride];

Same pairs, reassigned: thread t does the work thread 2×stride×t did. Now the working threads are 0, 1, 2… packed together, so only 4, then 2, then 1 warp has lanes on, and shared-load requests drop from 94 to 24 per block. But a new counter lights up: thread t touches word 2×stride×t, and neighbours that far apart share banks: 2-way at stride 1, up to 8-way at the deeper levels (red in the stepper). Each request now takes several passes, and the total comes back to 94 passes per block, the same as v1 for a completely different reason. Expected evidence: l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum jumps from zero to about 70 per block, while shared-memory wavefronts stay where v1 had them. Fewer requests, same work for the banks: a textbook “improvement” with a flaw that only the profiler shows.

v3: sequential addressing

for (unsigned stride = 128; stride > 0; stride >>= 1) {
    if (tid < stride) s[tid] += s[tid + stride];
    __syncthreads();
}

Fold the upper half onto the lower half. Active threads are packed (few warps send requests) and thread t touches words t and t+stride (consecutive, no conflicts): 24 requests and 24 passes per block, the minimum for this tree. What remains is structural, and identical in all three versions: every 256 input values pay for 8 barriers and 88 instructions per warp, for just 1 KB of data loaded per block. So the honest prediction for the benchmark is that v1, v2 and v3 land close together, v3 slightly ahead, nothing like the 2× steps of 2007. If your numbers show that, the lesson is the chapter’s whole point: optimization folklore is a hypothesis, and the SASS and the counters decide.

v4: coarsening

T sum = 0;                                            // each thread: 8 values, in a register
for (int k = 0; k < 8; ++k) {
    size_t idx = blockIdx.x * 256 * 8 + tid + k * 256;  // coalesced at every k
    if (idx < n) sum += in[idx];
}
s[tid] = sum;  __syncthreads();  /* ...then v3's tree */

Each thread first adds 8 values privately: no barrier, no shared memory, every thread busy, and 8 independent loads in flight. The block covers 2,048 values, so there are 8× fewer blocks and 8× fewer trees, barriers and atomics per byte. Expected evidence: DRAM throughput jumps toward the ceiling, barrier stalls shrink, long_scoreboard (waiting for DRAM) becomes the dominant stall, and total instructions drop sharply. This is usually the biggest single step in the sequence.

v5: warp shuffles for the last five levels

__device__ float warp_sum(float v) {
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xffffffff, v, offset);   // lane i reads lane i+offset's v
    return v;                                           // lane 0 holds the warp's sum
}

Once 32 or fewer values remain, they all live in one warp, and the lanes of a warp can exchange registers directly with __shfl_down_sync: no shared memory, no barrier. (The mask argument is chapter 2’s Volta lesson: it states which lanes take part instead of assuming lockstep.) v5 runs the shared-memory tree down to 64 values, then lets warp 0 finish with five shuffles. Five of eight barriers disappear, and so do five rounds of shared-memory loads and stores. Expected evidence: fewer shared-memory wavefronts and fewer barrier stalls, and probably almost no change in time, because v4 was already close to the DRAM ceiling. That is chapter 4’s lesson seen from the profiler: once you are on the roofline, removing work you do not wait for buys nothing. It would matter if the data were in L2, or if this reduction were fused into a kernel with other work, as it will be in Part II’s softmax.

v6: one wave, wide loads

The modern shape. Launch only as many blocks as fit on the GPU at once (48 SMs × 6 blocks of 256 threads = 288), let each thread walk the whole array with a grid-stride loop reading 16 bytes per load (float4), then reduce with shuffles within each warp, one shared-memory slot per warp, and shuffles again. 288 atomics in total instead of 32,768. It should sit at the ceiling, with long_scoreboard the main stall and the fewest instructions of any version. It is also the shape of reduction you will write inside fused kernels later.

The ceiling: CUB

CUB, NVIDIA’s library of parallel primitives (part of CCCL, shipped with the toolkit), has cub::DeviceReduce::Sum, tuned per architecture. The lab runs it as the reference, and ncu shows you how it does it: typically two kernels, one that reduces tiles to partial sums and a single-block kernel that finishes. In production you call the library; you write your own when the reduction must live inside another kernel.

7. The transposes, under ncu

Chapter 3 predicted why the three transposes differ: uncoalesced writes in the naive one, a 32-way bank conflict in the unpadded tile, neither in the padded one. make ncu-transpose collects exactly the counters that test those claims, on a 4096×4096 matrix, each kernel launched once. Here is what the mechanisms predict (numbers per warp-wide request, from the sector and bank rules):

Where the big numbers come from: 4096×4096 = 16.8 M floats are each read once out of the tile, 32 per warp-wide request, so there are 524,288 shared-memory load requests. Conflict-free, each is one wavefront (one pass through the banks). With 32 lanes on one bank, each needs 32 passes: 31 of them are counted as conflicts. The naive store needs 32 sectors per request because each lane writes a different row of the output, 16 KB from its neighbour (chapter 3’s scattered pattern).

Read your real numbers against this table. Small deviations are normal (a few requests at tile edges, the counters’ exact definitions). Large ones mean a claim in chapter 3 was wrong for your card, which would be worth knowing, and worth telling me.

8. The workflow

1. Timeline first. Run the real workload under nsys. Is the GPU busy? If there are gaps, you are overhead-bound (chapter 1), and no kernel work will help until they are gone: fuse, batch, use graphs, move CPU work off the critical path.
2. Pick the kernel with the largest total time in cuda_gpu_kern_sum. Amdahl decides where effort pays.
3. Classify it with Speed of Light and chapter 4’s roofline: compute-bound, memory-bound (which memory?), or neither.
4. Explain it: stall reasons, Memory Workload Analysis, the Source page. Find the line and the mechanism.
5. Write down one hypothesis that predicts a number: “padding removes the conflicts, so shared wavefronts per request drop from 32 to 1 and time drops by about X.”
6. Change one thing.
7. Re-measure both: the benchmark (did it get faster?) and the metric (did it change for the reason you thought?). A speedup for the wrong reason is a coincidence you cannot build on.

And keep a log: kernel version, change, hypothesis, benchmark result, key metrics. It takes a minute per iteration, it is the raw material of the write-ups that will become your portfolio, and it is exactly what an interviewer means by “walk me through how you optimized something.”

9. Lab 4: profile it yourself

cd ~/unibo-lessons/mlsys/labs/lab04-profiling
make test          # every reduction on ints (exact) and floats (tolerance); graph demo
make bench         # GB/s of v1..v6 and CUB; microseconds per kernel, three ways
make nsys          # timeline of the launch-gap demo + nvtx / API / kernel tables
make ncu-table     # this chapter's key metrics for every reduction, as text
make ncu-transpose # lab 2's transposes: sectors/request, wavefronts, conflicts
make ncu           # full report of every reduction: open it in the Windows GUI
make ncu KERNEL=reduce_v3   # ...or just one version

Check your understanding

Sources and further reading

• Hwu, Kirk, El Hajj, Programming Massively Parallel Processors, 4th ed. (2023), ch. 10: reduction trees, control and memory divergence, shared-memory staging, segmented reduction with atomics, thread coarsening.
• Mark Harris, Optimizing Parallel Reduction in CUDA (NVIDIA, 2007): the original seven-step version of this sequence. The numbers are from 2007 hardware; the reasoning is not dated.
• NVIDIA, Nsight Compute Profiling Guide: kernel replay, clock and cache control, metric naming, warp stall reasons (the descriptions in the stall widget are paraphrased from it). Also the Nsight Compute 2026.3 release notes (the --clock-control boost default since 2026.1; WSL requirements).
• NVIDIA, Nsight Systems User Guide: --trace, --cuda-graph-trace, nsys stats reports. Nsight Systems 2025.x release notes: the WSL timestamp known issue and its workaround.
• NVIDIA, CUDA on WSL User Guide: supported driver and Windows versions for developer tools.
• NVIDIA, CUDA C++ Programming Guide, sections on CUDA Graphs and warp shuffle functions; CUB documentation for cub::DeviceReduce.
• All metric names in this chapter were checked with ncu --query-metrics --chip gb205 (Nsight Compute 2026.3).