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

Building a roofline for your GPU

Day 30 published a roofline with an invalid row. The naive matmul read 2466 GB/s, 1008.5 percent of the same GPU's measured copy ceiling. No memory system moves ten times its own bandwidth.

The calculation used algorithmic bytes instead of bytes that reached DRAM. Only a hardware counter can provide the second value. This lesson rebuilds the roofline with dram__bytes.sum from Nsight Compute beside the hand-counted column and puts ten course kernels on it.

The difference between the columns measures cache effects.

The bytes you counted and the bytes that moved

A roofline needs three inputs: the GPU's bandwidth ceiling, its FMA ceiling, and one arithmetic intensity per kernel. Day 30 measured the two ceilings directly with a coalesced copy and a register-only FMA chain, and got the third input from the kernel source: count the flops, count the bytes the algorithm asks global memory for, divide.

That byte count has a name, compulsory traffic. It describes the algorithm, not the measured run.

The hardware keeps a different count. dram__bytes.sum is the number of 32-byte sectors that crossed the DRAM interface, multiplied by 32. An NVIDIA engineer describes these byte metrics as "not derived from any algorithmic property of your kernel, but measured in HW", and lists what that includes: "writing back from cache to DRAM, re-reading from DRAM, excessive sectors transferred due to uncoalesced accesses" (https://forums.developer.nvidia.com/t/how-to-compute-dram-bytes-read-sum-dram-bytes-read-sum/307220 , checked 2026-08-30).

Nothing in that list appears in a hand count, and the hand count's repeated reads that L1 and L2 served appear nowhere in the counter.

The two columns disagree on almost every kernel. Divide the measured bytes by the requested bytes to get one ratio per kernel. A ratio below one means the caches served traffic that the model assigned to DRAM.

A ratio above one means the hardware moved bytes the model never counted, because a 32-byte sector was fetched or written back to deliver 4 useful bytes. A ratio near one means the requested and measured traffic are close.

The ratio differs by kernel

The measured count is not always a smaller, corrected version of the hand count. Four pairs in the table show ratios in both directions.

The coalesced and stride-32 vector adds ask for identical bytes, 201 MB each, because day 11's permutation visits the same indices in a different order, breaking coalescing without changing the algorithm. The model puts them at the same point. The counter should separate them in the direction the model cannot express: the strided row moving more than it asked for, since every 4-byte float costs its warp a full sector and the L2 has evicted that sector long before the seven other threads that share it come round.

The naive and tiled matmuls ask for counts sixteen times apart, and both should produce nearly the same measured figure because three 512-square matrices use 3 MiB and the measured GPU's L2 is 4 MiB. Almost none of either kernel's traffic reaches DRAM. This explains day 30's 1008 percent row with a counter value.

The stencil pair tests day 34's cache claim. Day 34 measured tiling making its stencil slower, 0.283 ms against 0.173, and argued from the timer that the reads the tile removed were already cache hits. That argument was an inference.

If it is right, heatStepGlobal and heatStepTiled will show roughly equal DRAM traffic despite asking for very different amounts, and the shared-memory apron will have removed requests that the cache already served.

One program, two byte columns

Full program in code/day49-roofline/roofline.cu, with merge_dram.py beside it. The program and script depend on four choices.

Both ceilings use day 30's measured kernels unchanged, so the two lessons' ridge points are comparable. No vendor figure appears anywhere. A ceiling that the test cannot reach would add the same unknown error to every percentage.

    const double copyGBs = 2.0 * static_cast<double>(kElems) * sizeof(float) /
                           (static_cast<double>(copyMs) * 1.0e-3) / 1.0e9;
    const double fmaFlops = static_cast<double>(fmaThreads) *
                            static_cast<double>(kFmaIters) * kFmaChains * 2.0;
    const double peakGFlops =
        fmaFlops / (static_cast<double>(fmaMs) * 1.0e-3) / 1.0e9;
    const double ridge = peakGFlops / copyGBs;

Times come from events, bytes come from the profiler, and never the other way round. Nsight Compute replays a kernel and pins clocks with --clock-control base, so any rate it reports is a base-clock rate; a byte count does not move with the clock, which is why it is the one thing this lesson takes from the report. The program's --profile mode launches each kernel exactly once, no checks, no timing loops, so the ncu pass sees twelve launches instead of the 168 the normal path issues. Every launch configuration is written once, as a lambda used by both paths, so the profiled kernel cannot drift into a different shape from the timed one.

The asked-for column is printed from constants worked from the kernel source, so a wrong count is visible on the page rather than folded silently into a ratio:

    const double moved[kNumRows] = {3.0 * n * sizeof(float),
                                    3.0 * n * sizeof(float),
                                    2.0 * n * sizeof(float),
                                    2.0 * n * sizeof(float),
                                    (2.0 * d3 + d2) * sizeof(float),
                                    (2.0 * d3 / kTileDim + d2) * sizeof(float),
                                    (n + blocks) * sizeof(float),
                                    6.0 * interior * sizeof(float),
                                    (apron + interior) * sizeof(float),
                                    2.0 * n * sizeof(float)};

One new kernel fills the compute-bound part of the chart. Every kernel you have written so far sits left of the ridge point, so the flat half of the roofline would be empty. polynomialHorner runs a 256-step Horner chain over each loaded float, 513 flops for 8 bytes, which is the one intensity in the program high enough to be compute bound:

__global__ void polynomialHorner(const float* __restrict__ in,
                                 float* __restrict__ out, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        const float x = in[i] * kPolyScale;
        float acc = x;
#pragma unroll 16
        for (int k = 0; k < kPolyDegree; ++k) {
            acc = fmaf(acc, x, kPolyAdd);
        }
        out[i] = acc;
    }
}

One limit from day 30 remains: the FLOP counts are supplied by hand and no tool checks them. The program does what it can, a host recurrence that fails if the compiler shortens either FMA loop, but a miscounted numerator remains the easiest way to draw a wrong roofline.

Results

The timed half was re-verified without material drift on a Tesla T4 with driver 580.173.02 and CUDA 13.0 V13.0.88 on 2026-09-02. Its copy ceiling was 247.7 GB/s, FMA ceiling 7,489.8 GFLOP/s and ridge 30.24 FLOP/byte, consistent with the CUDA 12.6 profiler-backed record below. No CUDA 13 NCU pass was captured, so the existing 12.6 merged counters remain the 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; DRAM bytes from Nsight Compute 2024.3.2 under sudo at --clock-control base, merged against the timed run by merge_dram.py. Captured 2026-09-01 on the project's verification node; the transcript and the merged table are in the page's evidence file and the report ships at content/profiles/roofline/.

The merged table, which is the artifact this day exists to produce:

copy ceiling 245.1 GB/s, ridge point 30.49 FLOP/byte

kernel                   ask MB    L1 MB    L2 MB  DRAM MB DRAM/ask  %ceil
---------------------- -------- -------- -------- -------- -------- ------
vector add                201.3    201.3    201.5    219.9   1.0922  116.9
vector add, stride 32     201.3   1610.6   1610.9   4339.9  21.5566  107.6
transpose, naive          134.2    604.0    604.4    360.3   2.6844  111.5
transpose, tiled          134.2    134.2    134.5    159.3   1.1871  102.4
matmul, naive            1074.8    537.9     67.3      4.2   0.0039    3.8
matmul, tiled 16           68.2     67.8     65.1      3.9   0.0579    5.7
block reduction            67.4     69.2     69.6     92.2   1.3688   65.5
stencil, global           100.5    121.1     46.9     44.7   0.4448  112.0
stencil, tiled             39.0     48.5     47.0     48.2   1.2350  106.9
polynomial, 256           134.2    134.2    134.6    170.8   1.2724   60.7

One note before the results: the byte columns come from a cold profiler pass with the caches flushed between replays, so streaming rows include a few percent of measurement overhead and several are just over 100 percent of a ceiling measured warm. The conclusions use rows whose ratios differ from one by much more than that overhead.

Prediction one held everywhere it could. The naive matmul's DRAM traffic is 4.2 MB against a 1,074.8 MB ask: DRAM/ask 0.0039, an order of magnitude under the 0.03 bound, and the corrected percent-of-ceiling is 3.8 against day 30's published 1008.5. That row is now settled by a counter instead of an argument: the reads day 30 charged to DRAM were served by L2, and the L2 column's 67.3 MB says exactly where they went.

Prediction two split. The coalesced vector add came in at 1.0922, inside the 15 percent bound. The block reduction came in at 1.3688, outside it: 92.2 MB moved for a 67.4 MB ask, the atomics and partial traffic the ask column never counted plus the cold-capture overhead. The bound was too tight for any kernel that writes something other than a stream.

Prediction three failed. The strided add's DRAM/ask is 21.6, far above the predicted 4 to 8. The prediction counted the 32-sectors-per-request geometry and stopped there, but a sector miss does not move a sector, it moves DRAM's own larger granule, and a write-allocate stream adds more traffic.

Day 11's 8x described L1TEX traffic; DRAM moved more. The measured number replaces the predicted range.

Prediction four held, and day 34 is confirmed from counters. The two stencils moved 44.7 and 48.2 MB of DRAM, 7.8 percent apart, despite asks 2.6x apart. The reads tiling removed were cache hits, which is what day 34 could only infer from a timer.

Prediction five missed by 1.6 FLOP per byte. The ridge point came out at 30.49, under the predicted 32 to 34. The ceiling kernels are byte-identical to day 30's, but a measured ceiling is a measurement: this run's FMA ceiling measured 7,179 GFLOP/s against the copy ceiling's 245.1 GB/s, and their quotient sets the ridge for that run. The consequence predicted alongside still follows: against measured DRAM bytes the naive matmul's intensity is about 64 FLOP per byte, right of the measured ridge, so the kernel day 30 called memory bound is compute bound at this cache-resident size, and the two byte columns are two different true answers to "which ceiling binds".

The naive transpose had no prediction. It moved 360.3 MB for a 134.2 MB ask, giving DRAM/ask 2.68. L2 served some of the strided writes and passed the rest to DRAM.

The L2 column measured 604.4 MB for this row.

The result falls between 1.0 and the full sector waste, which shows why this row needed a profiler measurement.

Run it yourself

The timed half runs on a supported CUDA GPU:

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

The profiled half may need more permission. A stock driver can limit ncu to an administrator when it sets RmProfilingAdminOnly: 1.

Without access, you get 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. On a machine you administer, sudo ncu works.

Day 42 has the permanent modprobe option. The measured node had the parameter set to

  1. Current free-Colab behavior remains untested.

The shipped artifacts at content/profiles/roofline/ are the path: day49-roofline.ncu-rep, day49-bytes.csv, day49-bytes-warm.csv and day49-roofline.details.txt. merge_dram.py is standard library Python, needs no GPU, and produces the merged table from your own --json run or from the shipped one:

python3 merge_dram.py --json run.jsonl --ncu profile/day49-bytes.csv

The README carries the exact sudo ncu capture commands and what each flag is for. If you have no GPU at all, read /setup/learn-cuda-without-a-gpu: the merge step still works, because reading a report is not profiling.

Exercise

Before you open the merged table, write down a direction for each of the ten rows: DRAM/ask above 1, near 1, or below 1. Then produce the table, from your own GPU or from the shipped CSV, and score yourself. Explain the two rows you got most wrong.

Time: 30 to 45 minutes. Submit: ten directions, your score, the two explanations, and your GPU's ridge point if you ran the program.

Check: this is a measurement exercise, so the relations between rows must hold, not the absolute numbers. merge_dram.py exits non-zero if any kernel in the JSON has no dram__bytes.sum row, because a missing row is the difference between "the caches absorbed it" and "the profiler never saw that launch". The program itself refuses to print a roofline from a run whose kernels are wrong: fifteen checks, each a real branch returning EXIT_FAILURE, including a host recurrence that fails if the compiler shortened either ceiling's loop.

Hint 1

Two mechanisms separate the columns. Cache reuse lowers the measured DRAM count, while uncounted sector traffic raises it. For each kernel, ask which mechanism it has, and whether the working set fits in 4 MiB of L2.

Hint 2

The DRAM metric counts 32-byte sectors, while the hand count uses 4-byte floats. A kernel that uses all eight floats from each fetched sector has a ratio near 1.

Repeated cache hits can move the ratio below 1. Refetching a sector for each float can move it above 1.

Solution

Expect a ratio near 1 for the coalesced add, the reduction and the polynomial, the pure streams. Below 1 for both matmuls, massively, and for both stencils, moderately, because those four have reuse the caches capture.

Expect above 1 for the strided add, which wastes sectors without enough reuse. The naive transpose may fall on either side because it has store waste but also reuse of the written lines close enough in time that L2 may hide it, and the measured row is the answer.

The requested byte count describes the algorithm, while dram__bytes.sum describes the measured run. Their ratio shows cache and sector effects. Day 30's roofline was not wrong to count compulsory bytes; it was wrong to divide them by time and call the result bandwidth.

Pitfalls

ncu prints no metrics and one error. 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, and the shipped profile/ files where you have neither.

You took a rate from the profiler and a time from the program and compared them. --clock-control base pins the clocks during collection, so every rate ncu reports is a base-clock rate and the event timings ran at boost. Byte counts do not move with the clock; rates do. Take only counts from the report.

Your byte column and your time column came from different cache states. ncu flushes all GPU caches before each replay pass by default (--cache-control all), while the timed mean is ten back-to-back launches running warm. The day49-bytes-warm.csv capture uses --cache-control none --replay-mode application for the other state, and rows that move between the two files are telling you their traffic depends on what the previous launch left in L2.

You read a DRAM/ask of 0.003 as a broken join. It is the answer. A cache-resident matmul really does move a three-hundredth of what it asks for, which is why merge_dram.py prints that ratio to four decimals rather than two.

You trusted the ask column because it is printed by a program. It is arithmetic on the source, exactly as capable of being wrong as day 30's was, and no counter checks the FLOP numerator. You must check the values used on both roofline axes.

Go deeper

Next

Day 50 ends the module: three mystery kernels, their ncu reports, and no source code until you have committed to a diagnosis. It uses the metrics from this module, and the first task plots each kernel on the roofline.