The CUDA performance checklist, in diagnostic order

Three optimisations in this course made a kernel slower. Day 34 put a stencil in shared memory, cut its global reads from five per cell to 1.328, and went from 0.173 ms to 0.283. Day 17 capped a kernel's registers with __launch_bounds__, restored occupancy from 75 percent to 100, and went from 0.164 ms to 0.227.

Day 10 took the block size the occupancy API recommends and measured it slower than a block size a quarter as large.

Each change followed a common optimization rule, and in each case one number would have said so before the edit. This page is the order to ask for those numbers in, and the three kernels at the end are where you use it on code you did not write.

Five questions, and the one number that answers each

A checklist works best when each answer rules out as much work as possible. These five follow that order, not the order the tools present.

1. Is the kernel the slow part at all? Run Nsight Systems, not Nsight Compute, and read one number: nsys stats --report cuda_gpu_kern_sum totalled against the wall clock.

Day 9 is the whole argument for putting this first. Its kernel took 0.786 ms, the copies around it made the round trip 45.822 ms, and cudaSetDevice(0) alone cost 255.966 ms before any arithmetic ran.

A host clock around the launch read 0.016 ms. It measured only the queuing, which is launch overhead. If your kernels are under half the wall time, every step below is the wrong page.

2. Which ceiling is it under? Nsight Compute's Speed Of Light section, two rows: Compute (SM) Throughput [%] and Memory Throughput [%]. Whichever is higher names the limit you are near.

You can also get the answer without a profiler by dividing arithmetic intensity, flops over bytes, against your card's roofline ridge point, which day 30 measured at 33.09 FLOP per byte. Below it a kernel is memory bound whatever its arithmetic looks like; above it, compute bound.

This step decides which of the next two you run and which you skip.

3. If it is memory bound, is the traffic real?

Divide l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum by l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum. Four is the floor for floats, because a warp's 32 lanes on consecutive addresses span 128 bytes, which is four 32-byte sectors. Thirty-two is the ceiling, one sector per lane.

Day 11 measured that curve directly: requests stayed at 2,097,152 for every access pattern while sectors went 8,388,608, 16,777,216, 33,554,432 and then pinned at 67,108,864. The pattern cannot change how many instructions issue, only what each one costs, and coalescing is the difference between the two numbers.

4. If it is compute bound, are the lanes doing work?

smsp__thread_inst_executed_per_inst_executed.ratio. Thirty-two is the ceiling. Sixteen means half of every issue slot is masked off, which is what divergence looks like from the scheduler's side.

Day 22 priced it: a predicate on the lane index cost 2.02 times a warp-uniform one, and a predicate written as (threadIdx.x & 31) < 16 cost 2.00, the same branch written another way.

5. Is occupancy the limit, and would fixing it help? sm__warps_active.avg.pct_of_peak_sustained_active is achieved occupancy; the Occupancy section prints it next to Theoretical Occupancy [%] and next to the block-limit rows that name what caps it, registers or shared memory or block size.

Occupancy helps hide latency. If step 2 says you are already near a limit, more resident warps add nothing. If neither limit is near and occupancy is high, the answer is the dominant row of Warp State Statistics, which is where stall reasons live, and day 45 is the day that takes it apart.

Diagram: five questions, and the number that answers each. Five stacked bands read top to bottom, each with the question on the left, a single labelled bar or ratio in the middle, and the decision it forces on the right. Band 1, is the kernel the slow part: a wall-clock bar with a thin kernel segment inside it. Caption "0.786 ms of kernel inside a 45.822 ms round trip: 1.7 percent." Band 2, which ceiling: a point on a horizontal intensity axis with the ridge point marked. Caption "0.25 against a ridge point of 33.09 FLOP per byte: memory." Band 3, is the traffic real: 32 lane boxes over a row of sectors. Caption "4 sectors per request at the floor, 32 at the ceiling." Band 4, are the lanes working: 32 lane boxes, 16 of them hatched. Caption "16 threads per instruction, which day 22 measured as 2.02 times the time." Band 5, is occupancy the limit: two occupancy bars with their times beside them. Caption "100 percent and 0.227 ms against 75 percent and 0.164." Alt text: "Five questions in order, each answered by one number. Kernel time against wall time, intensity against a ridge point of 33.09, sectors per request, lanes per instruction, achieved occupancy against its time."

Why the order is not the order the tool gives you

Open a report and the first thing you meet is Speed Of Light, then the memory chart, then occupancy, then the warp state histogram. Read them in that order and you will spend an hour on a kernel that owns two percent of your run time. The tool is laid out by subject, and a diagnosis is not a subject.

Most people optimise by trying the thing they most recently read about, then measuring. That tests edits from a long list: coalescing, block size, occupancy, fusion, vector loads, spilling, and fast math. The five questions above test causes instead.

Each step is also cheaper than the one below it. Step 1 needs no counters at all, so it works on Colab, and it is the only step that can tell you the other four are irrelevant.

Step 2 halves what is left. Steps 3 and 4 are one division each. Step 5 is last because people often try it first, though it is the one that is least often the answer: day 17 raised a kernel to 100 percent occupancy and made it slower, and day 10 found nine of thirteen block sizes within five percent of the best.

There is a sixth question, and it is not on the list because it has no metric: is this the right algorithm? Day 30 draws that as the difference between moving a kernel up towards its ceiling and moving it right to a different one. The checklist only moves you up.

Three kernels that fail three different steps

Full program in code/day50-checklist/mystery.cu. Three kernels, four launches, and three ground rules.

The names carry no information, and that is deliberate. Every other file in this course names a kernel for what it does to memory. This one is the documented exception, because the exercise is to reach a diagnosis from the report. The source is in the repository and reading it will not tell you which step catches which kernel, which matches normal profiling: you almost always have the source and it almost never settles it.

Two of the three differ in exactly one thing. mysteryB and mysteryC run the same 1024 fused multiply-adds per thread and move the same one float in and one float out, so the branch is the only variable between them. mysteryC is then launched twice, once asking for 48000 bytes of dynamic shared memory it never reads and once asking for none, so occupancy is the only variable between those two.

Every launch is checked before anything is timed. Four correctness launches run first, each against a CPU reference at all 4,194,304 elements, with the output buffer poisoned beforehand so an element a kernel forgets to write returns a failure. That ordering also makes the profiler command simple: the first four launches ncu matches are one of each configuration.

mysteryA walks a 2048 by 2048 matrix with the row taken from the low bits of the thread index, so a warp's 32 addresses are 2048 floats apart:

__global__ void mysteryA(const float* __restrict__ in, float* __restrict__ out,
                         size_t n) {
    const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (t < n) {
        const size_t row = t & (kDim - 1);
        const size_t col = t >> kDimLog2;
        const size_t j = row * kDim + col;
        out[j] = fmaf(in[j], kScale, kBias);
    }
}

mysteryB and mysteryC share a device helper that runs four independent chains of fused multiply-adds. Four chains rather than one is what lets a single thread keep the arithmetic pipe busy without help from resident warps, and it is the reason the two mysteryC launches are a fair test of what occupancy is worth here rather than a rigged one:

__global__ void mysteryB(const float* __restrict__ in, float* __restrict__ out,
                         size_t n, int iters) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const float x = (i < n) ? in[i] : 0.0f;

    float y = 0.0f;
    if ((threadIdx.x & 1u) == 0u) {
        y = runChains(x, kMulA, kAddA, iters);
    } else {
        y = runChains(x, kMulB, kAddB, iters);
    }

    if (i < n) {
        out[i] = y;
    }
}

__global__ void mysteryC(const float* __restrict__ in, float* __restrict__ out,
                         size_t n, int iters) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    const float x = (i < n) ? in[i] : 0.0f;
    const float y = runChains(x, kMulA, kAddA, iters);

    if (i < n) {
        out[i] = y;
    }
}

The program answers step 5 without a profiler, using cudaFuncGetAttributes and cudaOccupancyMaxActiveBlocksPerMultiprocessor, which are ordinary runtime calls that need no counters and no root. Those two give registers per thread, shared memory per block, blocks per SM and theoretical occupancy on any card.

They cannot give achieved occupancy, because that is a measurement of what the scheduler did rather than of what it was allowed to do, and only the report has it.

Results

Re-verified on a Tesla T4 with driver 580.173.02 and CUDA 13.0 V13.0.88 on 2026-09-02. Launch attributes and every diagnostic conclusion reproduced. Absolute kernel times rose by roughly 2.4x, while the divergent/uniform ratio moved from 1.81x to 1.98x and the low-occupancy penalty stayed at 18 percent against the prior 19 percent.

The existing CUDA 12.6 NCU report remains the counter artifact of record.

Originally measured on a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo, profiled with Nsight Compute 2024.3.2 under sudo, --set full --launch-count 4 --clock-control base. Captured 2026-09-01 on the project's verification node; the transcript is in the page's evidence file and the report of record ships with its two text exports at content/profiles/performance-checklist/.

Part 1: what the launch configuration alone decides
kernel                      regs   static B   dynamic B  blk/SM warps/SM    occ %
-------------------------- ----- ---------- ----------- ------- -------- --------
mysteryA                       8          0           0       4       32    100.0
mysteryB                      14          0           0       4       32    100.0
mysteryC, 48000 B dynamic     14          0       48000       1        8     25.0
mysteryC, 0 B dynamic         14          0           0       4       32    100.0

Part 2: mean of 10 runs after 3 warm-ups, copies not included
kernel                      FLOP/byte    time ms      GB/s    GFLOP/s
-------------------------- ---------- ---------- --------- ----------
mysteryA                        0.250      0.449      74.7       18.7
mysteryB                      256.000      2.145      15.6     4004.9
mysteryC, 48000 B dynamic     256.000      1.414      23.7     6073.3
mysteryC, 0 B dynamic         256.000      1.184      28.3     7256.5

And the three report numbers the checklist steps ask for, one line per profiled launch:

kernel	sectors per request, ld / st	threads per executed instruction	achieved occupancy
mysteryA	31.91 / 32.00	32	81.90%
mysteryB	4.00 / 4.00	16.21	98.18%
mysteryC, 48000 B	4.00 / 4.00	32	24.82%
mysteryC, 0 B	4.00 / 4.00	32	86.29%

The four predictions, against the card:

One held. mysteryA came back at 31.91 load sectors per request and 32.00 on the store side, the other three at exactly 4. Lanes 2048 floats apart pay a full sector per lane; the chain kernels read one contiguous float per lane and pay one sector per eight.

Two held. mysteryB reads 16.21 threads per executed instruction, the lane-parity branch cutting the warp's useful width almost exactly in half, and it took 2.145 ms against 1.184 for mysteryC with the same occupancy, a ratio of 1.81, on the "about twice" the prediction priced from day 22.

Three died, and step 5 gets its worked example the other way. The two mysteryC launches are 1.414 and 1.184 ms, so the low-occupancy launch is 19 percent slower, past the tenth the prediction allowed. At 256 FLOP per byte this kernel is compute bound, and a compute-bound kernel at 25 percent occupancy has fewer warps ready to issue in the arithmetic pipes than the SM can consume; four chains per thread hide memory latency, not issue-slot starvation.

Day 45's within-2-percent pair moved 192 MiB and computed almost nothing, this pair computes 1024 FMAs per element and barely touches memory, and occupancy matters at the compute-heavy end that the checklist reaches last. The checklist order survives its own counterexample: steps 1 to 4 still clear both launches, and step 5 is where the answer was.

Four held. Both mysteryC launches are coalesced at 4 sectors per request, warp-uniform at 32 threads per instruction, and the 0-byte launch runs at 99.24 percent compute throughput, on its roof. For that launch the checklist ends by saying stop reading counters; the 48000-byte launch is the one the checklist catches at its last step.

Run it yourself

You do not need a GPU today and you do not need a profiler. That is the point of this day: the report ships with the lesson, the exercise is done from it, and min_cc is any for the first time since day 30.

If you have a CUDA GPU, run the program:

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o mystery mystery.cu
./mystery

If RmProfilingAdminOnly is 1, capturing counters needs an account with profiling permission:

sudo ncu --set full --kernel-name regex:'mystery' \
    --launch-count 4 --clock-control base \
    --export profile/day50-mystery ./mystery

Exercise

Open the shipped report and diagnose mysteryB and both launches of mysteryC. For each one, name the step of the checklist that catches it, the metric that settles it, and the single change you would make.

Time: 30 to 45 minutes. Submit: three verdicts, each with a metric name spelled the way the report spells it and the value you read off it.

Check: three questions marked in the browser from content/quizzes/day50.toml, with an explanation on every option, because the wrong options are where the teaching is. The program itself carries seven run-time checks, each a real branch returning EXIT_FAILURE rather than an assert, since CI builds Release and an assert under NDEBUG is deleted.

Four compare a launch against a CPU reference at every element and name the index and both values on failure. One rejects a zero time, because all four are about to become denominators.

One rejects the case where removing the shared memory request restores no occupancy, which is the assumption the whole occupancy half of this page rests on. The last checks that four rows printed, so this page and the program cannot drift on how many there are.

Hint 1

Two of the three launches are the same kernel. Before you read any counter, write down what can differ between two launches of one compiled kernel with identical arguments, and which of the five questions is the only one that can see that difference.

Hint 2

For mysteryB, step 2 will point you at the compute roof and step 4 has one number in it. Thirty-two is the ceiling. What fraction of the ceiling would you expect from a predicate that is true for even lanes and false for odd ones, and what does that fraction predict about the time?

Solution

mysteryB: step 4, and the metric is smsp__thread_inst_executed_per_inst_executed.ratio. Steps 1 to 3 clear it: the accesses are contiguous, so sectors per request is at its floor of 4, and the intensity is far to the right of the ridge point, so the memory chart has nothing to say. The ratio near 16 out of a possible 32 is the finding, and it means every warp is executing both sides one after the other while half its lanes sit masked.

The change is not always "rewrite the predicate". Here the two paths compute different chains, so threadIdx.x / 32 & 1 would give a different answer, not a faster one. The real move is to reorder the work so that a warp's 32 lanes want the same thing.

That is what day 37's warp-per-row kernel does to a matrix with uneven rows. Divergence is a property of how work is grouped, not of the if statement.

mysteryC with 48000 bytes: step 5, and the metric is achieved occupancy next to the Occupancy section's block-limit rows. Those rows name dynamic shared memory as what caps it at one block per SM. The change is to delete one launch argument, and on this card it bought 19 percent: 1.414 ms down to 1.184.

That number is the Results section's killed prediction three, and it is why the occupancy step exists at all, even last. Four chains per thread hide memory latency fine at 25 percent occupancy; they do not provide the extra ready warps a compute-bound kernel wants in its issue slots.

mysteryC with no dynamic shared memory: nothing catches it. Every step passes, which is a real answer and not a failure of the checklist. A kernel at its compute roof with coalesced access and no divergence is finished as written. Making it faster means changing what it computes, not how, and day 30 is the page that separates those two.

A checklist is a way of buying information in the order that rules out the most edits per question, and its best outcome is telling you to stop. Four of these five steps end in "then skip the next one". The one that never ends in an edit is the one people run first.

Pitfalls

You opened Nsight Compute first. It profiles kernels, so it cannot tell you that your kernels are two percent of your run time. Start with Nsight Systems, annotate the phases with NVTX ranges the way day 41 does, and only then pick a kernel.

ncu refuses to run. The message is "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", and the short form people search for is ERR_NVGPUCTRPERM. The cause is the driver parameter RmProfilingAdminOnly, which is 1 on this node. Use sudo ncu where you have root. This evidence does not settle current free-Colab behavior; every lesson ships its report so access is not required.

You compared a duration from the report against your own timing. Nsight Compute serialises kernels and replays each one many times to collect every counter, so a duration in a report is not what your program does when nobody is watching. Take counters from the report and times from CUDA events.

You pasted a metric name out of an old blog post. warp_execution_ efficiency, gld_efficiency and the rest are nvprof names and no longer exist. ncu --query-metrics lists what your card actually has, and the names on this page are the ones in the supplied reports.

You built without -lineinfo. The report still collects everything, but it cannot attribute a single counter to a source line, so the one section that would have pointed at the offending load is empty. It costs nothing at run time.

You read a percentage over 100 as a broken measurement. An analytical ceiling is built from the bytes the algorithm asks global memory for, and the caches serve part of them. Day 30 published a row at 1008 percent of its own copy ceiling for exactly that reason. In a report, dram__bytes.sum is what actually crossed the bus.

Go deeper

• Nsight Compute Kernel Profiling Guide, for the section list, the metric naming scheme and what each rule checks: https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html (checked 2026-08-30)
• Nsight Compute CLI reference, for --set, --kernel-name, --launch-count, --clock-control and the --import --page details export used here: https://docs.nvidia.com/nsight-compute/NsightComputeCli/index.html (checked 2026-08-30)
• Nsight Systems User Guide, for nsys stats and the cuda_gpu_kern_sum and cuda_api_sum reports that answer step 1: https://docs.nvidia.com/nsight-systems/UserGuide/index.html (checked 2026-08-30)
• NVIDIA's own page on the counter permission failure, with the full message and the modprobe option that clears it: https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters (checked 2026-08-30)
• Programming Massively Parallel Processors, 4th edition, chapter 6, on performance considerations and the order to apply them in: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0

Next

Module 6 opens on day 51 with CUDA streams, which is where step 1 stops being a diagnosis and starts being a design. Every kernel from here on is timed inside a pipeline rather than on its own, and the first question this page asks, whether the kernel is the slow part, becomes the thing you are building against.