100 DAYS OF GPU Course preview
Day 42Module 5
in-technical-review

Reading an Nsight Compute report

ncu may return no metrics when the driver limits counter access. It reports ERR_NVGPUCTRPERM, in full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device.

When the NVIDIA driver has RmProfilingAdminOnly: 1, only an administrator can use performance counters. On a machine you administer, run sudo ncu or change the driver setting.

This page ships two Nsight Compute reports and their text exports at content/profiles/nsight-compute/. They capture day 16's naive and tiled matmul on a Tesla T4.

You can complete the lesson by reading those files. The CUDA 12.6 privileged capture and the CUDA 13 ordinary-user permission failure are both listed in evidence.

The three numbers, and the order they answer in

A full report has a dozen sections and several hundred metrics. Three of them decide what you do next, and the rest are for after you know which question you are asking.

One: speed of light. The GPU Speed Of Light Throughput section appears first. Start with two rows.

Compute (SM) Throughput is what fraction of the arithmetic pipelines' peak rate the kernel reached. Memory Throughput is the same for the busiest memory unit. A kernel high on the first and low on the second is compute-bound; the other way round is memory-bound; low on both means it is stalled on latency and neither ceiling is the problem, which is the case warp stall reasons exist to explain.

The section also prints DRAM Throughput and Duration as separate rows, and the distinction between Memory Throughput and DRAM Throughput is the single most misread thing in this report. Older Nsight Compute labels those two headline rows Compute (SM) [%] and Memory [%], so a report captured a few years ago reads differently from the one shipped here; a 2022 transcript carrying the old spellings is at https://stackoverflow.com/questions/74245226/nsight-compute-not-showing-achieved-occupancy-in-the-metrics (checked 2026-08-30).

Two: the memory chart. Memory Workload Analysis is where the number turns into a cause. In the graphical report it is a diagram of boxes and arrows running from the kernel through the L1 and L2 caches out to global memory, each arrow labelled with the traffic that crossed it. On the command line the same numbers arrive as rows: Memory Throughput [Gbyte/s], L1/TEX Hit Rate, L2 Hit Rate.

The two counters that matter most here are the ones day 11 already used, l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum for how many warp-wide load instructions issued and l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum for how many 32-byte sectors those instructions pulled. Their ratio is sectors per transaction. Compare that ratio and the total sector count across the two kernels.

Three: achieved occupancy. The Occupancy section prints Theoretical Occupancy and Achieved Occupancy next to each other, plus the four block limits (Block Limit SM, Block Limit Registers, Block Limit Shared Mem, Block Limit Warps) that produced the theoretical figure. Theoretical occupancy is arithmetic on the launch configuration and you can get it before the kernel runs; achieved occupancy is the average number of resident warps over the kernel's life and needs a run.

When they differ, the launch could have filled the machine and the workload did not, and barriers, an uneven tail and one block taking longer than its neighbours are the usual reasons. Read it third, because it explains a gap rather than naming a bound.

Diagram: where a load instruction goes. Three horizontal bands. Each has a warp on the left, then a box for shared memory, a box for L1, a box for L2 and a box for DRAM to the right, with arrows carrying the traffic between them. Band 1, the naive kernel: every arrow from the warp runs past shared memory into L1 and on to L2. Caption "1,024 global loads per thread, two sectors each, 16,777,216 sectors across the grid." Band 2, the tiled kernel's global half: a thin bundle of arrows on the same path. Caption "64 global loads per thread, four sectors each, 2,097,152 sectors across the grid." Band 3, the tiled kernel's shared half: a thick bundle of arrows that stops at the shared memory box and never reaches L1 or L2. Caption "1,024 shared loads per thread, and not one of them fetches a sector." Alt text: "Tiling moves load instructions rather than removing them. The naive kernel issues 1,024 global loads per thread. The tiled kernel issues 64 global loads and 1,024 shared ones, and a shared load fetches no sector."

Why the naive matmul is not a coalescing problem

After day 11 the reflex on seeing a slow kernel is to suspect the access pattern, and here the reflex is wrong. Both matmuls are already coalesced.

The block is 16 by 16 with threadIdx.x fastest, so a warp of 32 lanes is two rows of sixteen consecutive columns. Look at what that does to the naive kernel's two loads:

acc += a[row * k + p] * b[p * n + col];

a[row * k + p] depends only on row, so the warp asks for two addresses, one per half, and they land in two sectors. b[p * n + col] depends only on col, so the warp asks for sixteen consecutive floats, which is 64 bytes starting on a 64-byte boundary, and that is two sectors as well. Two sectors per request is close to ideal.

Day 11's worst case was thirty-two.

The tiled kernel's global loads are worse on that ratio, four sectors per request rather than two, because each of the warp's two rows now reads a run of sixteen instead of sharing one. It is still eight times cheaper overall, because it issues sixteen times fewer of them.

The ratio alone does not show the cost. The naive kernel has the better sectors-per-request figure and moves eight times the sectors, and no amount of fixing the access pattern would have found that, because there was nothing wrong with the access pattern.

Capturing a report someone else can read

Full program in code/day42-ncu/matmul_profile.cu. The kernels are day 16's, unchanged. Four rules make the capture reproducible.

One size, and it divides the tile. 512 by 512 by 512 with a 16-wide tile means no bounds guard ever fires, so the report describes the steady-state shape rather than an edge case. Day 16 owns the ragged sizes and the guards they need. 512 cubed is also day 30's roofline row, which is what lets this report answer a question that day left open.

A fixed launch order, so the capture command can name a launch. Each kernel launches fourteen times here: one correctness launch, three warm-ups, ten timed. The reports profile the first timed launch, which is why the capture passes --launch-skip 4.

That option counts only launches matching the kernel filter, so with one kernel named per run the four skipped are that kernel's own. The program prints the command with the number filled in from its own constants, so the README and the code cannot drift apart about which launch was captured.

The prediction is printed before the counters are read. These constants say what the memory rows should hold, and the derivation is in the comment above them:

constexpr size_t kSectorBytes = 32;
constexpr size_t kBlocks = (kDim / kTileDim) * (kDim / kTileDim);
constexpr size_t kWarps = kBlocks * ((kTileDim * kTileDim) / kWarpSize);
constexpr size_t kNaiveRequests = kWarps * 2 * kDim;
constexpr size_t kTiledRequests = kWarps * 2 * (kDim / kTileDim);
constexpr size_t kNaiveSectorsPerRequest = 2;
constexpr size_t kTiledSectorsPerRequest = 4;
constexpr size_t kNaiveSectors = kNaiveRequests * kNaiveSectorsPerRequest;
constexpr size_t kTiledSectors = kTiledRequests * kTiledSectorsPerRequest;

constexpr size_t kNaiveGlobalLoads = 2 * kDim;
constexpr size_t kTiledGlobalLoads = 2 * (kDim / kTileDim);
constexpr size_t kTiledSharedLoads = 2 * kTileDim * (kDim / kTileDim);

// A, B and C read or written once each. This is the floor on DRAM traffic,
// not a prediction of it: both kernels re-read A and B many times and the
// caches decide how much of that reaches memory.
constexpr size_t kCompulsoryBytes = 3 * kDim * kDim * sizeof(float);

The last three lines support this claim: tiling does not cut how many load instructions a thread issues. It moves them. Two loads per k step is two loads per k step either way.

The naive kernel issues 1,024 global loads per thread, while the tiled kernel issues 64 global and 1,024 shared ones. The tiled kernel issues more load instructions and is faster, because a shared-memory load fetches no sector from L2. A static_assert ties the tiled kernel's shared-load count to the naive kernel's global-load count, counted two different ways, because this page's whole reading of the memory chart rests on them being equal.

Timings come from a plain run, counters from a profiled one. Nsight Compute serializes the kernel it is collecting and replays it once per pass, so a cudaEvent figure captured under ncu measures the profiler. Never take both from the same run.

Note. On Turing the L1 data cache and shared memory are one 96 KB unit per SM, carved either 64 KB shared and 32 KB L1 or the other way round (https://docs.nvidia.com/cuda/turing-tuning-guide/index.html , checked 2026-08-30). The tiled kernel's shared reads therefore land in the same hardware its global reads left. Whether L1/TEX Cache Throughput falls, holds or rises is not obvious from the source, and this capture settles it.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), profiled with Nsight Compute 2024.3.2 under sudo, --clock-control base. Captured 2026-09-01 on the project's verification node; the transcript is in the page's evidence file and both reports ship with their text exports at content/profiles/nsight-compute/.

CUDA 13.0 (V13.0.88) re-verification on 2026-09-02, with driver 580.173.02 on the same Tesla T4, passed both kernels against the CPU reference. The ordinary-user ncu attempt was blocked by ERR_NVGPUCTRPERM with RmProfilingAdminOnly: 1, while the new Nsight Systems capture completed. The counter table below remains the CUDA 12.6 privileged capture; no CUDA 13 counter values are claimed.

Rows and metric names are spelled as the report spells them, so a reader can search for either. The one derived line says so.

Report row or metric naive tiled
Compute (SM) Throughput 60.78% 72.67%
Memory Throughput 60.78% 72.67%
DRAM Throughput 1.25% 1.95%
Duration 1.18 ms 744.64 us
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 8,388,608 524,288
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum 16,777,077 2,083,878
sectors per request, derived 2.0 4.0
dram__bytes_read.sum + dram__bytes_write.sum 4,741,440 B 4,641,024 B
L1/TEX Hit Rate 87.58% 5.13%
L2 Hit Rate 97.93% 97.05%
Theoretical Occupancy 100% 100%
Achieved Occupancy 95.29% 94.94%

The four predictions, against the card:

The request counts matched the derivation exactly and the sector counts came within 0.7 percent. 8,388,608 and 524,288 requests, to the digit. The sector counters read 16,777,077 against a derived 16,777,216 and 2,083,878 against 2,097,152, both a hair under, which is the counters seeing a handful of sectors served before they started counting rather than an error in the geometry. Sectors per request lands on 2.0 and 4.0 as derived.

Neither kernel is anywhere near DRAM bound, and day 30's flaw is settled. Total DRAM traffic is 4.74 MB for the naive kernel and 4.64 MB for the tiled one, within 1.6x of the 3,145,728-byte compulsory floor and about 220 times below the 1,074.8 MB day 30 charged the naive kernel when it computed bytes by hand. The row that read over a thousand percent of the copy ceiling was counting reads the caches served: the report says so directly, with DRAM Throughput under 2 percent for both kernels while Memory Throughput sits over 60.

The memory chart separates the kernels and the duration does not. Duration differs by 1.59x, sectors moved by 8.05x. A reader handed only the times could not say why; a reader handed the sector counts can. The L1/TEX Hit Rate row is the same story from the other side: the naive kernel hits L1 on 87.58 percent of its sectors because a warp's row re-reads are cache-resident, and the tiled kernel's global loads hit L1 on 5.13 percent because shared memory already absorbed the reuse and what is left is compulsory.

Achieved occupancy sits below theoretical for both, and lower for the tiled kernel by less than half a point. 95.29 against 94.94 percent. The direction held, the margin is nothing: two barriers per tile step cost this kernel almost no residency at this shape, and the prediction's reasoning bought more than the card charged.

Do not memorise the percentages. The useful parts across GPUs are the order you read the sections in and the fact that two of these rows answer different questions with the same word in them.

Run it yourself

Reading the shipped reports needs nothing: ncu --import works without counters, because reading a report is not profiling. The published copies live at content/profiles/nsight-compute/; the verification run captures them into code/day42-ncu/profile/ and the build copies them over.

Capturing them needs performance-counter permission. On a system you administer, use sudo ncu or the persistent driver setting in the README.

If you cannot collect counters, read the shipped .ncu-rep files. The .nsys-rep in the same directory records an Nsight Systems trace that does not need those counters. See /setup/learn-cuda-without-a-gpu if you do not have a GPU.

Exercise

Four questions, all answerable from this page and the shipped reports.

Time: 25 to 40 minutes. Submit: four answers, each naming the report row or the counter that settles it.

Check: the quiz in content/quizzes/day42.toml, marked in the browser, with an explanation on every option. The four questions above have their answers in the reveal below, and each one points at something you can look up in profile/day42-naive.details.txt or profile/day42-tiled.details.txt rather than at a feeling about how the hardware behaves. An answer that names no row and no counter is not finished.

  1. Read the naive kernel's Memory Throughput row and its DRAM Throughput row. Day 30 called the same kernel bandwidth-bound. Are those two rows and day 30 talking about the same bandwidth?
  2. The tiled kernel issues more load instructions per thread than the naive one and runs faster. Which row of the memory chart explains that, and which row cannot?
  3. Suppose the two kernels report the same Theoretical Occupancy and different Achieved Occupancy. Name two things on this page that could cause that gap and one that could not.
  4. You are an ordinary user on a system where RmProfilingAdminOnly is 1. Which of the shipped captures can you make, and which is blocked?
Hint 1

Three questions ask you to distinguish two similar metrics. Identify which metric measures the hardware unit in question.

Hint 2

For question 2, count load instructions per thread for each kernel and then count sectors. The two counts move in opposite directions.

Solution

1. No. Memory Throughput in the speed-of-light section is the busiest memory unit as a fraction of its own peak, and on this kernel that unit is L1TEX, not DRAM. DRAM Throughput is a separate row and it is the one that means what day 30 meant.

The obvious wrong answer is that both say "memory-bound", which is true of the word and false of the bottleneck.

Fixing DRAM traffic on a kernel that is saturating L1 buys nothing.

2. The sector count explains it; the request count cannot. The naive kernel issues 1,024 global loads per thread and the tiled kernel 1,088 loads in total, 64 global and 1,024 shared. Counting instructions, the tiled kernel does more work.

Counting sectors pulled through L1TEX from L2, it does an eighth as much, because a shared-memory load fetches no sector. The obvious wrong answer is "it issues sixteen times fewer loads", which is true only of the global ones and skips where the other 1,024 went.

3. Barriers and an uneven tail could; registers and shared memory could not. The tiled kernel waits at two __syncthreads() per tile step, and the last wave of blocks leaves part of the machine idle.

Registers and shared memory are exactly what Block Limit Registers and Block Limit Shared Mem already priced into the theoretical figure, so if theoretical occupancy is equal they are not what separated the achieved figures. The obvious wrong answer is "the tiled kernel uses shared memory, so fewer blocks fit", which the block-limit rows refute in the same section.

4. The .nsys-rep, not the .ncu-rep files. Nsight Systems tracing does not need these performance counters, while Nsight Compute does. With RmProfilingAdminOnly set to 1, an ordinary user gets ERR_NVGPUCTRPERM.

The report gives you sectors, hit rates, and occupancy. Your model of the kernel explains those values. Here, knowing that a warp covers two rows of sixteen columns tells you what to change.

Pitfalls

ncu prints no metrics and one error. The string is ERR_NVGPUCTRPERM, in full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device. Performance counters are admin-only under the stock driver default RmProfilingAdminOnly: 1. Use sudo ncu where you have root, set NVreg_RestrictProfilingToAdminUsers=0 where you own the machine, and read the shipped report where you have neither.

You timed the kernel under the profiler. Nsight Compute replays the kernel once per collection pass, so every clock in the program is measuring the profiler. Run bare for times, run under ncu for counters, and say on the page which run each number came from.

You read Memory Throughput as DRAM bandwidth. It is the busiest memory unit against its own peak, and DRAM Throughput is the row you wanted. A kernel can sit near the top of the first and near the bottom of the second, which is exactly what a cache-resident working set looks like.

You profiled the first launch of a kernel. Lazy module loading is the default from CUDA 12.2 on Linux, so the first launch of each kernel pays its own load, and the report of that launch describes a cost you will never pay again. Skip past the warm-ups. Day 9 makes the same point about a timer.

You compared two reports captured at different clocks. --clock-control base pins the GPU clocks for the collection, and without it a card that was warm for one capture and cold for the other gives you two rooflines rather than one. Both captures here pass it.

You optimised for the occupancy row because the report prints it. Occupancy buys latency hiding and nothing else, and past the point where the latency is covered the registers it costs buy less. Day 45 tunes one kernel two ways to equal speed at different occupancies.

Go deeper

Next

Day 43 changes this kernel in three steps. It grades each step with report ratios rather than absolute times, which vary by GPU.

Day 45 studies the occupancy row. Day 49 rebuilds day 30's roofline from dram__bytes.sum instead of a byte estimate, then computes memory bandwidth and arithmetic intensity from measured traffic.