100 DAYS OF GPU Course preview
Day 30Module 3
in-technical-review

Arithmetic intensity and the memory-bound mindset

Someone on r/CUDA, after a week of optimisation work:

"I assumed the kernel was compute-bound because the mathematical operations seemed complex ... the problem was actually quite simple: it was memory-bound due to non-coalesced global memory accesses, so those fancy changes were completely useless."

https://www.reddit.com/r/CUDA/comments/1qkp3lb/my_first_optimization_lesson_was_stop_guessing_lol/ (checked 2026-08-29)

The work was not wrong. It was aimed at the wrong ceiling, and one division would have said so before any of it started. Count the floating point operations your kernel does, count the bytes it asks memory for, divide.

That ratio decides which limit you are under. You can get it from the source before the profiler or the first run.

This is a checkpoint, so nothing new is introduced. You take kernels you have already written, place five of them on one chart, and find out which optimisations can possibly help each one.

One division decides which ceiling you are under

A GPU has two speed limits and your kernel is under exactly one of them.

The first is memory bandwidth: bytes per second between the SMs and DRAM. The second is arithmetic throughput: floating point operations per second, which for FP32 is the fused multiply-add rate across every lane on the card. Neither is a property of your kernel.

Both are properties of the machine. You can measure them in about a second.

What your kernel supplies is one number, its arithmetic intensity:

intensity = FLOP performed / bytes moved

Put those three numbers together and you get the roofline. At intensity I, the fastest your kernel can possibly run is

attainable = min(peak FLOP/s, peak bytes/s * I)

The two terms cross at one point, and that point has a name:

ridge point = peak FLOP/s / peak bytes/s

Below the ridge point the second term is smaller, so bandwidth decides your speed and no amount of cleverer arithmetic changes it. Above it, the arithmetic decides and no amount of cleverer memory access changes it. The vocabulary is Nsight Compute's own: it calls the sloped part the memory bandwidth boundary, the flat part the peak performance boundary, and the meeting point the ridge point (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html section 2.9, "Roofline Charts", checked 2026-08-30).

The ridge point is the one number worth carrying out of this page, because it is a fact about your card and it does not change when you rewrite a kernel. A kernel is memory bound or compute bound by comparing one ratio against it.

Four kernels, worked by hand

None of this needs a GPU. Four kernels you have already written, counted from their source.

Vector add, day 5. The expression is out[i] = a[i] + b[i]. It performs one add and moves three floats across the bus, two in and one out, for 12 bytes.

1 FLOP / 12 bytes = 0.083 FLOP per byte

The lowest intensity in this course, and the first kernel every beginner writes. Day 9 measured it 24 times faster than one CPU thread and the whole program 2.4 times slower. This ratio explains the result: the kernel does too little arithmetic to offset its memory traffic.

Naive transpose, day 12. out[col * n + row] = in[row * n + col]. Four bytes in, four out, and not one arithmetic operation on the way.

0 FLOP / 8 bytes = 0 FLOP per byte

Zero, not "a small number", so on a log axis it has no position at all. The roofline has one thing to say about a transpose and day 12 spends its whole page on it: the sloped ceiling is the only one it will ever meet, so coalescing is the entire optimisation surface.

Block reduction, day 24. Each element is read once, four bytes, and contributes one add. A 256-thread block does 255 adds and writes one float back.

1 FLOP / 4 bytes = 0.25 FLOP per byte

That is the ceiling for a streaming kernel. One operation per element, each element read exactly once, is the best case for anything that walks an array, and it is still nowhere near a ridge point.

Tiled matmul, day 16. Worth doing slowly, because it is the only kernel here whose intensity you can change.

One output element needs K multiply-adds, and an FMA is two flops, so 2K flops. Without tiling each step loads one element of A and one of B, which is 8K bytes:

2K FLOP / 8K bytes = 0.25 FLOP per byte

The same as the reduction, for a kernel that looks a hundred times heavier. Now tile it: with a T by T tile each thread loads two floats per tile step and there are K / T steps, so 8K / T bytes:

2K FLOP / (8K / T) bytes = T / 4 FLOP per byte

At day 16's T of 16 that is 4 flops per byte, sixteen times the naive kernel, from a change that touched no arithmetic at all. The general form is the tile dimension over the element size in bytes, so the same tile over FP16 gives 8.

The distinction matters. Coalescing, block size and occupancy move a kernel up towards the ceiling it is already under.

Tiling and fusion move it right, to a different ceiling. Only the second kind can help a kernel that is already at 90 percent of its bandwidth limit.

The formula also caps how far tiling alone can go. Reaching a ridge point of R needs T above 4R, and two T by T float tiles cost 8T^2 bytes of shared memory. An SM with 64 KiB runs out at T of 90.

Past a ridge point in the low twenties, a shared tile cannot get you there at all. That is why days 43 and 44 add register tiling. You can reach that conclusion without writing either kernel.

Where the count comes from, and where it lies

The intuition that sends people the wrong way is that a kernel with complicated-looking arithmetic must be compute bound. The forum post at the top of this page is that intuition, held by someone competent, costing a week.

It fails because the two are unrelated. sqrtf and expf look heavy in source and are a handful of SASS instructions. A float4 load looks like nothing and is 16 bytes off the far side of the bus.

The second trap is in the count itself. Bytes here means the traffic the algorithm asks global memory for, not the traffic DRAM sees, because the caches serve part of the request. The naive matmul is the clean case: every thread in a block reads the same row of A, the model charges every one of those reads to DRAM, and L1 serves most of them.

So a measured point can sit above the sloped ceiling the model puts it under. That is not a broken measurement. It means the caches did work the model did not know about, and the gap measures the reuse the cache supplied.

Nsight Compute draws a line per level of the hierarchy for that reason. Day 49 adds them.

One smaller caveat: a ceiling measured with a copy is half reads and half writes, so a read-only kernel can beat it slightly. It is a good ceiling, not a law.

Measuring the two ceilings instead of quoting them

Full program in code/day30-roofline/roofline.cu. It measures both ends of the roofline and then places five kernels on it.

Three rules hold it together, and the first is the one every published roofline breaks.

No vendor number appears in it. A roofline drawn against a figure off the box is a roofline nobody can reach, so every percentage on it is wrong by the same unknown factor.

Both ceilings come from a kernel that ran here. The bandwidth one is a coalesced copy.

It uses the CUDA C++ Best Practices Guide's own effective-bandwidth formula, ((Br + Bw) / 10^9) / time (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html section 9.2.2, "Effective Bandwidth Calculation", checked 2026-08-30). The other is a chain of fused multiply-adds with nothing but registers inside the loop.

    // The two ceilings. Both are measured. Neither is a vendor figure, and
    // the widget on the page refuses one for the same reason: a ceiling you
    // cannot reach tells you nothing about how close you got.
    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;

Every kernel runs in this one program. The five points are re-run rather than quoted from days 5, 12, 16 and 24. Borrowing four days' numbers would fold four buffer sizes, four block shapes and four thermal states into one chart, and the chart would still look tidy.

Everything comes off one card in one run. The program times each kernel with events after a warm-up.

The counts are printed beside the times. They are worked from the kernel source and nowhere else, so a wrong count is visible on the page rather than folded silently into a ratio.

    const double n = static_cast<double>(kElems);
    const double blocks = static_cast<double>(vecBlocks);
    const double d3 = static_cast<double>(kMatDim) *
                      static_cast<double>(kMatDim) *
                      static_cast<double>(kMatDim);
    const double d2 = static_cast<double>(matElems);

    const char* names[kNumRows] = {"vector add", "transpose, naive",
                                   "block reduction", "matmul, naive",
                                   "matmul, tiled 16"};
    const double flops[kNumRows] = {n, 0.0, n - blocks, 2.0 * d3, 2.0 * d3};
    const double moved[kNumRows] = {
        3.0 * n * sizeof(float), 2.0 * n * sizeof(float),
        (n + blocks) * sizeof(float), (2.0 * d3 + d2) * sizeof(float),
        (2.0 * d3 / kTileDim + d2) * sizeof(float)};
    const double milliseconds[kNumRows] = {vectorAddMs, transposeMs, reduceMs,
                                           matmulNaiveMs, matmulTiledMs};

The FP32 ceiling is the one kernel here you have not met. It carries eight independent chains per thread so an FMA never waits on the one before it, and takes its iteration count as a run-time argument, because a constexpr count would let nvcc unroll all 16,384 multiply-adds into straight-line code and measure the instruction cache instead:

__global__ void fmaPeak(float* __restrict__ out, int iters) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;

    float acc[kFmaChains];
#pragma unroll
    for (int q = 0; q < kFmaChains; ++q) {
        acc[q] = static_cast<float>(q) * 0.125f;
    }

    for (int k = 0; k < iters; ++k) {
#pragma unroll
        for (int q = 0; q < kFmaChains; ++q) {
            acc[q] = fmaf(acc[q], kFmaMul, kFmaAdd);
        }
    }

    float sum = 0.0f;
#pragma unroll
    for (int q = 0; q < kFmaChains; ++q) {
        sum += acc[q];
    }
    out[i] = sum;
}

One caveat. A compiler that decided that loop was dead would report a ceiling the card cannot reach, and every point would then be measured against it, so the program recomputes one thread's chain on the host and compares. Nothing else in it can detect a wrong count, which is the model's own warning: you supply the FLOP and byte figures and no tool checks them for you.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5), 40 SMs
16777216 floats per buffer (64 MiB), matmul 512 cubed, tile 16

The two ends of this card's roofline, both measured
  copy, coalesced, in and out            0.549 ms     244.5 GB/s
  fma, 8 chains, 2048 iterations         1.327 ms    8092.1 GFLOP/s
  ridge point                                         33.09 FLOP/byte
  Below that intensity nothing on this card can be compute
  bound, whatever its arithmetic looks like.

Five kernels from days 5, 12, 16 and 24, placed on it
kernel                FLOP/byte       ms     GB/s   GFLOP/s   bound by  %ceil
--------------------  ---------  -------  -------  --------  ---------  -----
vector add                0.083    0.786    256.2      21.3  bandwidth  104.8
transpose, naive          0.000    1.302    103.1       0.0  bandwidth   42.2
block reduction           0.248    0.563    119.7      29.7  bandwidth   48.9
matmul, naive             0.250    0.436   2466.1     615.9  bandwidth  1008.5
matmul, tiled 16          3.938    0.275    247.8     976.0  bandwidth  101.3

5 of the 5 sit left of the ridge point, so for those the sloped
ceiling is the one that binds and a faster inner loop buys nothing.
FLOP/byte counts the traffic the algorithm asks for. The caches
serve part of it, so a row can read above 100 percent of a ceiling
the model says it should be under.

The CUDA 13.0 re-run materially changed the measured compute ceiling. Copy bandwidth stayed close, 244.5 to 249.9 GB/s, but the eight-chain FMA kernel fell from 8,092.1 to 3,741.4 GFLOP/s. The resulting ridge moved from 33.09 to 14.97 FLOP/byte.

The transpose, reduction and both matrix multiplies also took roughly two times as long, while vector add was stable. All five intensities still sit left of the lower ridge, so every bound classification and the lesson's rule still holds.

The measured ridge point was 33.09 FLOP per byte with CUDA 12.6 and 14.97 with CUDA 13.0. Below the ridge measured in the same run, this card cannot be compute-bound because it runs out of bandwidth first. Every kernel in the first thirty days sits left of both measurements.

That is why almost every optimisation so far has moved fewer bytes rather than done less work. The ridge is a measured input to the model, not a fixed GPU specification.

One row is wrong on purpose to show you how to spot it. The naive matmul reports 2466 GB/s, which is 1008 percent of this card's measured copy ceiling. No memory system moves ten times its own bandwidth.

The program computes GB/s from compulsory bytes, the minimum the algorithm must read, and a naive matmul re-reads the same rows and columns constantly.

L1 and L2 serve those repeat reads before they reach DRAM. Dividing compulsory bytes by elapsed time therefore describes traffic that did not happen.

Any row above 100 percent of the ceiling shows this effect. The tiled matmul at 101.3 percent is the same effect, nearly gone, because tiling is exactly what removes the repeats.

This is the checkpoint's main lesson. A roofline built from an analytical byte count tells you what the algorithm demands, not what the hardware did. To get the second you need dram__bytes from Nsight Compute, which module 5 covers.

Until then, read the intensity column as a property of the algorithm. A percentage over 100 means the cache did work the model did not count.

Run it yourself

You do not need a GPU today. Every intensity above is arithmetic, the quiz is answerable from this page, and that is why this day's minimum compute capability is any.

If you have a CUDA GPU, run:

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

There is no Compiler Explorer embed here. The program allocates 192 MiB, launches seven kernels thirteen times each, and checks every one of them against a CPU reference, including a 512 cubed matmul and a 4096 square transpose.

That is more than a 20 second run cap supports. If you have no GPU, read /setup/learn-cuda-without-a-gpu.

Exercise

Three questions first, none of them needing a GPU. The method is on this page; the kernels are yours from earlier days. Work out the arithmetic intensity of each and say which ceiling it is under.

  1. The grayscale kernel from day 7.
  2. The shared-memory transpose from day 13.
  3. The 5 by 5 Gaussian blur from day 20, first as a kernel that reads every tap from global memory, then as the tiled one day 20 uses.

Then, if you have a card, run the program and write down your ridge point.

Time: 30 to 45 minutes. Submit: four ratios with the byte counts you used, four verdicts, and your card's ridge point.

Check: the three questions are marked in the browser, from content/quizzes/day30.toml, with an explanation on every option because the wrong options are the interesting part. The program checks itself in nine places, each a real branch returning EXIT_FAILURE rather than an assert, because CI builds Release and an assert under NDEBUG is deleted.

Seven compare a kernel against a CPU reference. One of those recomputes the FMA chain on the host, so it fails if the compiler shortened the loop and inflated the FP32 ceiling. The last two check that no timing is zero and that five rows printed.

Your ridge point should be a number in the tens. A value below 1 means the FMA kernel did not run.

Hint 1

Count what the code does, not what the formula means, and count the two things separately. For the bytes, ask what crosses the bus rather than what the source names: a value a thread reads twice is one trip if it stayed in a register or in shared memory, and two if it did not.

Hint 2

For the blur, each output pixel touches 25 input pixels and each input pixel is touched by 25 outputs. Only one of those two numbers belongs in your byte count, and which one it is depends on a decision day 13 taught you to make.

Solution

1. Grayscale, day 7: zero flops.

Day 7's luma is fixed point. It uses three integer multiplies, two adds and a shift, with no floating point operation in the kernel.

On a FLOP roofline its intensity is zero, like the transpose but for a different reason. The chart can say only that it is bandwidth bound. If you answered about 1.25 flops per byte, you counted the formula rather than the code.

A FLOP roofline cannot see integer work. Many kernels use integer work.

2. Shared-memory transpose, day 13: zero, memory bound. Tiling through shared memory does not change a transpose's intensity at all, because the numerator was zero and stays zero.

It changes how close the kernel gets to the sloped ceiling by replacing a strided write with a coalesced one.

This optimisation moves the point up, not right. Tiling does not always raise intensity.

3. The blur, day 20: 0.48 untiled and about 5.5 tiled, memory bound either way. Twenty-five taps is 25 fused multiply-adds, so 50 flops, and the filter resides in __constant__ memory, so it is not new traffic per pixel.

Charge every tap to global memory and the denominator is 25 reads plus one write, 104 bytes, giving 0.48.

Load the tile once and share it. The traffic is one read and one write per pixel, which would give 6.25, but day 20's 36 by 36 apron feeds a 32 by 32 tile, so the value is nearer 5.5. That one choice raises intensity by more than ten times without changing the arithmetic.

Then look at what day 20 does next. Its separable version computes the same blur in two 5-tap passes, 20 flops instead of 50, with the global traffic unchanged.

Its intensity falls to about 2.2, and it is faster. Intensity is a diagnosis, not a score to maximise.

One line to keep: an optimisation either moves a kernel up towards its ceiling or right to a higher one, and the ratio is what tells you which one you need. Coalescing, block size and occupancy are all "up". Tiling, fusion and keeping data resident are "right".

Pitfalls

Counting an FMA as one flop. It is two, a multiply and an add, and every convention that quotes GFLOP/s counts it that way. Halving your intensity can move a kernel across the ridge point on paper.

Using the number on the box. A vendor peak is a clock rate times a lane count and no kernel reaches it. Both of your ceilings should come from a kernel that ran.

The Best Practices Guide keeps these apart on purpose, with theoretical bandwidth in section 9.2.1 and effective bandwidth in 9.2.2 (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html , checked 2026-08-30).

Charging the caches to DRAM. The intensity you compute by hand is what the algorithm asks for. A kernel with reuse moves less than that, which is why a point can sit above the sloped ceiling.

Treat the gap as a measure of reuse, not a broken chart.

Optimising the arithmetic of a memory-bound kernel. The forum post at the top of this page. Fast math, cheaper transcendentals and a tighter inner loop all change the numerator, and the numerator is not what is limiting you.

Reading the roofline as a diagnosis. It says which ceiling applies, not why you are at 40 percent of it, and it cannot: occupancy, latency, divergence and access pattern all live in that gap. Day 42 reads the speed-of-light section that does answer it.

Go deeper

Next

Day 31 is the prefix sum, a pattern whose intensity is as low as vector add's and which tiling cannot fix, because every output depends on every input before it. The ratio you learned to compute today says in advance which ceiling it is under.