GPU Fundamentals: Profiling Metrics & Roofline

Optimizing a kernel starts with one question: what is it waiting on, memory or arithmetic? The roofline model answers it on paper, and the profilers confirm it. Every program below ran on a Tesla T4 with CUDA 13.0, and ncu/nsys are the Nsight 2025.3 tools that ship with it.

The two roofs of your GPU

A GPU has two speed limits: how many floating-point operations per second it can do, and how many bytes per second its memory can deliver. Both come from device attributes.

// peaks.cu
#include <cstdio>
#include <cuda_runtime.h>

int main() {
  int sms, clockKHz, memClockKHz, busBits;
  cudaDeviceGetAttribute(&sms, cudaDevAttrMultiProcessorCount, 0);
  cudaDeviceGetAttribute(&clockKHz, cudaDevAttrClockRate, 0);
  cudaDeviceGetAttribute(&memClockKHz, cudaDevAttrMemoryClockRate, 0);
  cudaDeviceGetAttribute(&busBits, cudaDevAttrGlobalMemoryBusWidth, 0);
  const int fp32PerSM = 64;                            // FP32 lanes per SM on Turing (7.5)
  double flops = 2.0 * sms * fp32PerSM * clockKHz * 1e3;   // an FMA counts as 2 FLOPs
  double bytes = 2.0 * memClockKHz * 1e3 * busBits / 8;    // two transfers per memory clock
  printf("SMs: %d, boost clock: %d MHz, memory clock: %d MHz, bus: %d bits\n",
         sms, clockKHz / 1000, memClockKHz / 1000, busBits);
  printf("peak FP32:      %.0f GFLOP/s\n", flops / 1e9);
  printf("peak bandwidth: %.0f GB/s\n", bytes / 1e9);
  printf("ridge point:    %.1f FLOPs per byte\n", flops / bytes);
  return 0;
}

$ nvcc -o peaks peaks.cu && ./peaks
SMs: 40, boost clock: 1590 MHz, memory clock: 5001 MHz, bus: 256 bits
peak FP32:      8141 GFLOP/s
peak bandwidth: 320 GB/s
ridge point:    25.4 FLOPs per byte

40 SMs x 64 FP32 lanes x 2 FLOPs per fused multiply-add x 1.59 GHz gives 8.1 TFLOP/s. The memory side is 256 bits x 2 transfers x 5 GHz, or 320 GB/s. Those are the T4’s published figures.

The ridge point divides one by the other. A kernel has to do 25 FLOPs for every byte it moves before the arithmetic, not the memory, becomes the limit. The FP32 lane count isn’t a device attribute. It’s fixed per architecture (64 for compute 7.5), so check NVIDIA’s table if you run this on another GPU.

Operational intensity on paper

A kernel’s operational intensity (OI) is its FLOPs divided by the bytes it moves. Its best possible speed is the lower of the two roofs: min(peak FLOP/s, OI x bandwidth).

# oi.py
PEAK_GFLOPS, PEAK_GBS = 8141, 320         # Tesla T4, from peaks.cu

kernels = {                               # (FLOPs, bytes moved) per element, float32
    "y = a * x":           (1, 8),
    "c = a + b":           (1, 12),
    "y = a * x + y":       (2, 12),
    "fma loop, 4096 x 4":  (2 * 4 * 4096, 8),
}

for name, (flops, nbytes) in kernels.items():
    oi = flops / nbytes
    roof = min(PEAK_GFLOPS, oi * PEAK_GBS)
    bound = "memory" if oi * PEAK_GBS < PEAK_GFLOPS else "compute"
    print(f"{name:20} OI {oi:8.3f} FLOP/B -> at most {roof:5.0f} GFLOP/s, {bound}-bound")

$ python3 oi.py
y = a * x            OI    0.125 FLOP/B -> at most    40 GFLOP/s, memory-bound
c = a + b            OI    0.083 FLOP/B -> at most    27 GFLOP/s, memory-bound
y = a * x + y        OI    0.167 FLOP/B -> at most    53 GFLOP/s, memory-bound
fma loop, 4096 x 4   OI 4096.000 FLOP/B -> at most  8141 GFLOP/s, compute-bound

c = a + b reads two floats and writes one (12 bytes) for one addition, an OI of 0.083. Even at full bandwidth it can’t top 27 GFLOP/s, a third of one percent of the T4’s arithmetic peak. All three elementwise kernels sit two orders of magnitude left of the 25.4 ridge.

The fourth kernel is the opposite case. It loads one float, does 4,096 rounds of 4 FMAs in registers, and stores one float, so memory barely matters. The next two programs measure both kinds.

Measuring a memory-bound kernel

Time c = a + b over 32M floats with CUDA events, then turn the time into bandwidth and FLOP/s and compare each with its roof.

// vecadd.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void add(const float* a, const float* b, float* c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) c[i] = a[i] + b[i];
}

int main() {
  const int n = 1 << 25;                              // 32M floats per array
  float *a, *b, *c;
  cudaMalloc(&a, n * 4); cudaMalloc(&b, n * 4); cudaMalloc(&c, n * 4);
  cudaEvent_t s, e; cudaEventCreate(&s); cudaEventCreate(&e);
  add<<<n / 256, 256>>>(a, b, c, n);                  // warm-up
  cudaEventRecord(s);
  add<<<n / 256, 256>>>(a, b, c, n);
  cudaEventRecord(e); cudaEventSynchronize(e);
  float ms; cudaEventElapsedTime(&ms, s, e);
  double gbs = 12.0 * n / ms / 1e6, gflops = 1.0 * n / ms / 1e6;
  printf("%.2f ms: %.0f GB/s (%.0f%% of 320), %.1f GFLOP/s (%.1f%% of 8141)\n",
         ms, gbs, 100 * gbs / 320, gflops, 100 * gflops / 8141);
  return 0;
}

$ nvcc -o vecadd vecadd.cu && ./vecadd
1.57 ms: 257 GB/s (80% of 320), 21.4 GFLOP/s (0.3% of 8141)

The model predicted at most 27 GFLOP/s; the kernel reached 21.4, 80% of the bandwidth roof. The arithmetic units sat almost idle at 0.3% of peak. Making the addition faster would change nothing, because every warp is waiting on memory.

Three runs gave 257 to 262 GB/s. In practice, 75-90% of the theoretical peak is about as good as a streaming kernel gets.

Measuring a compute-bound kernel

Now the far side of the ridge. Four independent FMA chains per thread keep the arithmetic pipeline full while memory traffic is close to nothing.

// fma.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void fma_loop(float* x, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  float a = x[i], b = a + 1, c = a + 2, d = a + 3;    // 4 independent chains
  for (int k = 0; k < iters; k++) {
    a = a * 0.999f + 0.001f; b = b * 0.999f + 0.001f;
    c = c * 0.999f + 0.001f; d = d * 0.999f + 0.001f;
  }
  x[i] = a + b + c + d;
}

int main() {
  const int n = 1 << 22, iters = 4096;
  float* x; cudaMalloc(&x, n * 4); cudaMemset(x, 0, n * 4);
  cudaEvent_t s, e; cudaEventCreate(&s); cudaEventCreate(&e);
  fma_loop<<<n / 256, 256>>>(x, iters);               // warm-up
  cudaEventRecord(s);
  fma_loop<<<n / 256, 256>>>(x, iters);
  cudaEventRecord(e); cudaEventSynchronize(e);
  float ms; cudaEventElapsedTime(&ms, s, e);
  double gflops = 2.0 * 4 * iters * n / ms / 1e6;     // 4 FMAs per pass, 2 FLOPs each
  printf("%.2f ms: %.0f GFLOP/s (%.0f%% of 8141)\n", ms, gflops, 100 * gflops / 8141);
  return 0;
}

$ nvcc -o fma fma.cu && ./fma
17.06 ms: 8054 GFLOP/s (99% of 8141)

8,054 GFLOP/s is 99% of the arithmetic roof. This kernel does about 375 times the FLOP rate of vecadd, on the same GPU, because its OI puts it under the compute roof.

The four chains matter. Each FMA depends on the previous result in its own chain, so one chain alone would leave the pipeline waiting on its own latency.

This kernel is the most sensitive to clock speed, because it does nothing but arithmetic. Across four runs it printed between 5648 GFLOP/s (69%) and 8054 GFLOP/s (99%). The T4 lowers its clock when it’s hot or has just woken up, and a compute-bound kernel slows down in step. Run it a few times, and treat the best steady result as the kernel’s speed.

Confirming it with Nsight Compute

ncu reads the GPU’s hardware counters for a kernel. Two metrics give the verdict straight away: DRAM throughput and SM throughput, each as a percentage of its peak. --launch-skip 1 skips the warm-up launch and profiles the timed one.

$ M=dram__throughput.avg.pct_of_peak_sustained_elapsed,sm__throughput.avg.pct_of_peak_sustained_elapsed
$ ncu --launch-skip 1 --launch-count 1 --metrics $M ./vecadd
==PROF== Connected to process 5496 (/content/gpu/vecadd)
==PROF== Profiling "add" - 0 (1/1): 0%....50%....100% - 6 passes
1963.45 ms: 0 GB/s (0% of 320), 0.0 GFLOP/s (0.0% of 8141)
==PROF== Disconnected from process 5496
[5496] vecadd@127.0.0.1
  add(const float *, const float *, float *, int) (131072, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    -------------------------------------------------- ----------- ------------
    Metric Name                                        Metric Unit Metric Value
    -------------------------------------------------- ----------- ------------
    dram__throughput.avg.pct_of_peak_sustained_elapsed           %        89.77
    sm__throughput.avg.pct_of_peak_sustained_elapsed             %        23.41
    -------------------------------------------------- ----------- ------------

$ ncu --launch-skip 1 --launch-count 1 --metrics $M ./fma
==PROF== Connected to process 5548 (/content/gpu/fma)
==PROF== Profiling "fma_loop(float *, int)" - 0 (1/1): 0%....50%....100% - 6 passes
2150.25 ms: 64 GFLOP/s (1% of 8141)
==PROF== Disconnected from process 5548
[5548] fma@127.0.0.1
  fma_loop(float *, int) (16384, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 7.5
    Section: Command line profiler metrics
    -------------------------------------------------- ----------- ------------
    Metric Name                                        Metric Unit Metric Value
    -------------------------------------------------- ----------- ------------
    dram__throughput.avg.pct_of_peak_sustained_elapsed           %         0.25
    sm__throughput.avg.pct_of_peak_sustained_elapsed             %        99.74
    -------------------------------------------------- ----------- ------------

The profiler agrees with the roofline. vecadd keeps DRAM 90% busy while the SMs sit at 23%: memory-bound. fma_loop is the mirror image, with SMs at 99.7% and DRAM at a quarter of a percent: compute-bound.

Repeated profiles agreed to within 0.05 percentage points. Ignore the program’s own timing lines under ncu: it printed 1963 ms and “0 GB/s”. ncu replays the kernel (“6 passes”) and saves and restores memory around each pass, so wall-clock timings measured inside the program are useless while it’s attached. Take times from normal runs, and metrics from ncu.

Don’t expect the two percentages to match: the same kernel shows 80% in vecadd’s own output and 90% here. ncu measures against the unit’s peak over the kernel’s own elapsed cycles, with the GPU clocks locked (its default --clock-control base). The event timing in vecadd includes launch overhead and runs at whatever clock the GPU picked. They measure different things.

Fusion: fewer bytes for the same work

Computing out = a * x + b with two kernels writes the intermediate y to DRAM and reads it back. A fused kernel keeps y in a register.

// fusion.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void scale(const float* x, const float* a, float* y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) y[i] = a[i] * x[i]; }
__global__ void add(const float* y, const float* b, float* out, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) out[i] = y[i] + b[i]; }
__global__ void fused(const float* x, const float* a, const float* b, float* out, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) out[i] = a[i] * x[i] + b[i]; }

int main() {
  const int n = 1 << 24, g = n / 256; float *x, *a, *b, *y, *out;
  for (float** p : {&x, &a, &b, &y, &out}) cudaMalloc(p, n * 4);
  cudaEvent_t e0, e1, e2; cudaEventCreate(&e0); cudaEventCreate(&e1); cudaEventCreate(&e2);
  fused<<<g, 256>>>(x, a, b, out, n);                 // warm-up
  cudaEventRecord(e0);
  scale<<<g, 256>>>(x, a, y, n); add<<<g, 256>>>(y, b, out, n);
  cudaEventRecord(e1);
  fused<<<g, 256>>>(x, a, b, out, n);
  cudaEventRecord(e2); cudaEventSynchronize(e2);
  float t1, t2; cudaEventElapsedTime(&t1, e0, e1); cudaEventElapsedTime(&t2, e1, e2);
  printf("two kernels: %.3f ms (6 array passes)\nfused:       %.3f ms (4 array passes)\n", t1, t2);
  return 0;
}

$ nvcc -o fusion fusion.cu && ./fusion
two kernels: 1.575 ms (6 array passes)
fused:       1.028 ms (4 array passes)

The two-kernel version reads x, a, writes y, then reads y, b and writes out: six array passes. The fused kernel reads x, a, b and writes out: four. Both are memory-bound, so time follows bytes, and 6/4 = 1.5 is almost exactly the 1.53x speedup measured. Three runs agreed to within 0.02 ms.

The FLOPs didn’t change at all. Fusion raised the OI from 2/24 to 2/16 FLOPs per byte by cutting traffic. For a memory-bound kernel, that’s the lever that works.

Seeing the timeline with Nsight Systems

nsys records every kernel launch, with no hardware counters needed. Its kernel summary is the quickest way to find where a program’s GPU time goes.

$ nsys profile -o fusion --force-overwrite true ./fusion > /dev/null 2>&1
$ nsys stats -q --force-export=true --report cuda_gpu_kern_sum fusion.nsys-rep | sed 's/ *$//'

 ** CUDA GPU Kernel Summary (cuda_gpu_kern_sum):

 Time (%)  Total Time (ns)  Instances  Avg (ns)   Med (ns)   Min (ns)  Max (ns)   StdDev (ns)                                Name
 --------  ---------------  ---------  ---------  ---------  --------  ---------  -----------  ----------------------------------------------------------------
     56.4        1,996,840          2  998,420.0  998,420.0   996,132  1,000,708      3,235.7  fused(const float *, const float *, const float *, float *, int)
     21.8          773,770          1  773,770.0  773,770.0   773,770    773,770          0.0  scale(const float *, const float *, float *, int)
     21.8          771,050          1  771,050.0  771,050.0   771,050    771,050          0.0  add(const float *, const float *, float *, int)

fused shows two instances, the warm-up and the timed launch, at about 0.998 ms each. scale and add take 0.77 ms each, 1.54 ms together: the same 1.5x gap the events measured. nsys doesn’t replay kernels, so these durations are trustworthy.

sed only strips the trailing spaces from the table. --force-export=true makes nsys stats rebuild its SQLite export when you re-profile into the same file name.

Common gotchas

• Optimizing the wrong roof. Fewer instructions won’t speed up a kernel at 90% DRAM throughput. Cut bytes: fuse, reuse, or use smaller types.
• Taking timings under ncu. It replays kernels; time with events or nsys instead.
• One chain of dependent math. A single FMA chain per thread can’t reach the compute roof; give the scheduler independent work.
• Comparing against the spec sheet only. About 80% of theoretical bandwidth is normal for streaming. ncu’s “pct_of_peak_sustained” percentages use a different baseline (locked clocks, the kernel’s own cycles), so don’t compare them directly with event-timed bandwidth.

Practice

1. Add a kernel y = x * x * x * x (one read, one write) to the dictionary in oi.py. What does the model say?

2. In fma.cu, try iters of 4, 16, 32, 64 and 256. Predict where the kernel stops being memory-bound, then compare.

3. Profile fusion with the same two ncu metrics. What do you expect for scale, add and fused?

4. nsys reported 2 instances of fused but 1 each of scale and add. Why?