Kernel optimization & the Roofline model.

Before you optimize, find out what is limiting you. The Roofline model answers one question — "is this kernel compute-bound or memory-bound?" — and that answer tells you which optimizations can possibly help. Then there is a small, well-studied toolbox of techniques (the optimization taxonomy from Hijma et al.): improve memory access, increase parallelism, reduce instruction overhead, use the memory hierarchy, and exploit hardware units. This article is the method: measure the ceiling, find the bottleneck, apply the matching technique, repeat.

Phase 1 gave you the rules the hardware rewards. But knowing the rules is not the same as knowing which rule to apply to this kernel. Spending a day perfecting coalescing on a kernel that is actually compute-bound is wasted effort. Optimization without measurement is guessing. So Phase 2 starts with the single most useful analytical tool in GPU performance — the Roofline model — and then organizes the techniques into a taxonomy you can match against the bottleneck it reveals.

Arithmetic intensity: the number that decides everything

Every kernel has a fundamental ratio: how much arithmetic it does per byte it moves from memory. This is arithmetic intensity (AI), measured in FLOPs per byte.

It is a property of the algorithm and its data movement, not of the hardware. Some examples make it concrete:

Operation	Arithmetic intensity	Character
Vector add (c = a + b)	~0.08 FLOP/byte	deeply memory-bound
SAXPY / element-wise scale	~0.17 FLOP/byte	memory-bound
Attention / softmax (naive)	~1–10 FLOP/byte	memory-bound
Dense matrix multiply (large, tiled)	~50–250 FLOP/byte	compute-bound
Convolution (well-tiled)	~tens FLOP/byte	borderline → compute-bound

Vector add does one add for every 12 bytes it touches (two 4-byte reads, one 4-byte write) — its intensity is hopelessly low, so it will always be limited by memory bandwidth no matter how fast the ALUs are. Matrix multiply, by contrast, does O(N³) work on O(N²) data, so its intensity grows with size and it becomes compute-bound. This single ratio predicts which wall you will hit.

The Roofline model

The Roofline model (Williams, Waterman & Patterson) plots achievable performance (FLOP/s, vertical) against arithmetic intensity (FLOP/byte, horizontal), both on log scales. It draws two ceilings — the "roofline":

• A sloped roof on the left: the memory-bandwidth ceiling. Performance here = arithmetic intensity × peak memory bandwidth. Low-intensity kernels are capped by how fast you can stream bytes.
• A flat roof on the right: the peak compute ceiling. Performance here = the GPU's peak FLOP/s. High-intensity kernels are capped by raw arithmetic throughput.

The two meet at the ridge point — the arithmetic intensity at which a kernel transitions from memory-bound to compute-bound. For an H100 with ~3.35 TB/s HBM3 and ~67 TFLOP/s FP64 (or far more for Tensor-Core formats), the ridge sits around 20–60 FLOP/byte depending on the precision you compare. Anything to the left of the ridge is memory-bound; anything to the right is compute-bound.

The Roofline model's power is diagnostic. Plot your kernel's measured performance as a point. If it sits far below the roofline directly above it, you have headroom — and the kernel's position relative to the ridge tells you which direction to push:

• Memory-bound and below the sloped roof: you are wasting bandwidth. Apply memory techniques — coalescing, caching/reuse in shared memory, fewer/larger transactions. Adding FLOPs will not help.
• Compute-bound and below the flat roof: you are wasting ALU throughput. Apply compute techniques — better instruction mix, Tensor Cores, lower precision, removing redundant work, fixing divergence.
• Already near the roofline: the kernel is close to optimal for its intensity. The only way up is to change the algorithm's intensity — e.g. fuse operations so you move fewer bytes per FLOP, pushing the point rightward toward the ridge.

The optimization taxonomy

Hijma et al.'s survey of GPU optimization techniques organizes the field into a handful of families. Rather than a grab-bag of tricks, think of them as categories you select from based on what the Roofline told you. The five families:

1. Memory access optimizations

The highest-leverage family for the (very common) memory-bound case. Covered in depth in article 1.2: coalescing global access, Struct-of-Arrays layout, using shared memory for reuse, avoiding bank conflicts, and using read-only/constant caches (__ldg, __constant__) for data all threads share. The goal is fewer bytes moved and every moved byte used.

2. Increasing parallelism & occupancy

Make sure the SMs have enough independent work to hide latency: right-size thread blocks, expose more parallelism in the algorithm, manage register/shared-memory pressure so occupancy is sufficient. This is where the Occupancy Calculator earns its keep (next section). Remember the lesson from 1.2: aim for enough occupancy, not maximum.

3. Reducing instruction overhead

Do less work per useful result: loop unrolling (#pragma unroll) to cut loop-control instructions, fast math intrinsics (__sinf, __expf, --use_fast_math) when precision allows, replacing expensive division/modulo with shifts and multiplies, and minimizing redundant address computation. Most useful for compute-bound kernels.

4. Exploiting the memory hierarchy & data reuse

The fusion/tiling family: tiling to maximize reuse, register blocking (each thread computes several outputs so loaded values are reused in registers), and kernel fusion to keep intermediates on-chip. This is what raises arithmetic intensity and moves you toward the ridge.

5. Exploiting specialized hardware

Use the units the GPU built for your problem: Tensor Cores for matrix math (via CUTLASS, cuBLAS, or WMMA intrinsics), asynchronous copy (cp.async, TMA) to overlap memory movement with compute, warp-level primitives (__shfl_sync for register-to-register exchange without shared memory), and lower precision (FP16/BF16/FP8) to double or quadruple effective throughput and halve memory traffic.

The Occupancy Calculator in practice

From article 1.2 you know occupancy is a min() over register, shared-memory, and block limits. The Occupancy Calculator turns that into concrete numbers. The modern workflow:

# 1. See per-thread register + shared-memory usage at compile time
nvcc -Xptxas -v -o kernel kernel.cu
# ptxas info: Used 48 registers, 12288 bytes smem

# 2. Query the achievable occupancy from the runtime API
int blocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, myKernel, 256, 0);
# blocks = how many 256-thread blocks fit per SM

# 3. Let the runtime pick a good block size for you
int minGrid, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGrid, &blockSize, myKernel, 0, 0);

The calculator's real value is showing you the limiting resource. If it reports that registers cap you at 50% occupancy, you can cap registers with __launch_bounds__(256, 4) (telling the compiler "I want at least 4 blocks of 256 threads per SM"), which trades a little per-thread register budget for more resident warps — and then you measure whether it actually helped, because, as always, more occupancy is not automatically faster.

Optimizing ML kernels specifically

The Niteesh & Ampareeshan analysis of optimizing ML models with CUDA highlights how these general techniques specialize for deep learning, where the workload is dominated by big matrix multiplies and element-wise ops:

• Mixed precision is the biggest single lever. Training and inference in FP16/BF16 (with FP32 accumulation) roughly doubles Tensor-Core throughput and halves memory traffic — helping on both axes of the Roofline. FP8 on Hopper pushes this further. This is why every modern training stack defaults to mixed precision.
• Fuse the element-wise tail. A matmul followed by bias-add, activation, and dropout should be one fused kernel, not four. Each separate kernel would round-trip the activations through HBM; fusion keeps them in registers/shared memory. Libraries like cuDNN and compilers like Triton/torch.compile do this automatically, but understanding why lets you spot when they fail to.
• Pick the right library before writing a kernel. cuBLAS and cuDNN contain hand-tuned, architecture-specific GEMM and convolution kernels that are extremely hard to beat. The optimization skill is often knowing when to call them vs when a fused custom kernel (CUTLASS, Triton) wins — typically when your shape is unusual or you can fuse surrounding ops the library cannot.
• Batch to raise intensity. A single-sequence inference is memory-bound (low intensity); batching many requests amortizes the weight loads over more compute, sliding the operation toward the compute roof. This is the core economics of inference serving and connects directly to Phase 3.

A worked optimization session

Here is the method end to end on a hypothetical slow kernel, in the order you would actually do it:

1. Measure first. Run the kernel under Nsight Compute (article 2.2). It reports achieved FLOP/s, memory throughput, and where you sit on the Roofline. Say it shows 15% of peak memory bandwidth and 3% of peak compute — clearly memory-bound and far below the sloped roof.
2. Check coalescing. The profiler's memory-transaction efficiency is 18%. That points straight at a strided access. Fix the indexing so thread i hits element i. Re-measure: bandwidth jumps to 60%.
3. Add reuse. The kernel re-reads the same input across threads. Tile it into shared memory. Now arithmetic intensity rises and the point moves right; bandwidth pressure drops and the kernel speeds up again.
4. Tune occupancy. The Occupancy Calculator says shared memory now limits you to 33%. Shrink the tile slightly or adjust block size; re-measure. Maybe it helps, maybe it does not — only the measurement decides.
5. Consider fusion / precision. If a neighboring element-wise op exists, fuse it. If precision allows, drop to FP16. Re-measure against the new ceiling.
6. Stop when near the roofline. Once the point is close to the ceiling above it, further micro-optimization has diminishing returns. The kernel is as fast as its algorithm allows on this hardware.

Common pitfalls

FAQ

How do I find my kernel's arithmetic intensity without hand-counting FLOPs?

Nsight Compute reports it directly — it counts executed floating-point instructions and DRAM bytes and can plot your kernel on a Roofline chart for you. You rarely need to compute it by hand; you use the hand calculation to sanity-check the profiler and to reason about algorithm changes.

Is the Roofline ridge point the same for every GPU?

No — it depends on the ratio of peak compute to peak bandwidth, which differs by GPU and by precision. An H100 comparing FP8 Tensor-Core peak against HBM3 has a very different ridge than the same chip comparing FP64. Always Roofline against the specific peak FLOP/s for the precision you actually use.

Does lower precision ever hurt?

Numerically, yes — FP16 has limited range and can overflow/underflow, which is why mixed precision keeps a FP32 master copy and accumulates in FP32. FP8 needs careful scaling. The speed is nearly free; the cost is the engineering to keep training stable, which mature frameworks now handle for you.

When is loop unrolling worth it?

When loop-control overhead is a meaningful fraction of the loop body — short, hot inner loops with small trip counts. For long loops with heavy bodies, unrolling barely helps and can hurt by raising register pressure and lowering occupancy. Let the profiler decide.

What's the difference between cuBLAS and CUTLASS?

cuBLAS is a closed, pre-tuned library you call. CUTLASS is an open template library that lets you build GEMM-like kernels and fuse custom epilogues into them. Use cuBLAS for standard matmuls; use CUTLASS when you need to fuse or customize. Both are covered in article 2.2.

Takeaways

• Arithmetic intensity (FLOP/byte) is a property of your algorithm and decides which wall you hit.
• The Roofline model diagnoses memory-bound vs compute-bound and shows your headroom to the ceiling.
• The optimization taxonomy — memory access, parallelism/occupancy, instruction overhead, hierarchy/reuse, specialized hardware — is selected by the Roofline diagnosis.
• Fusion raises intensity by removing HBM round-trips — the highest-leverage change for memory-bound ML kernels.
• Use the Occupancy Calculator to find the limiting resource; tune it, then measure.
• Mixed precision helps both axes and is the single biggest lever for deep learning kernels.
• Never optimize without measuring before and after.

Next, article 2.2 covers the toolchain that makes all of this measurable and writable: Nsight Systems and Nsight Compute for profiling, and CUTLASS, TensorRT, Triton, and Helion for actually producing fast kernels without hand-writing PTX.

References & further reading

• Williams, Waterman & Patterson — Roofline: An Insightful Visual Performance Model — the original Roofline paper.
• Hijma et al. — Optimization Techniques for GPU Programming — the survey behind the technique taxonomy.
• Niteesh & Ampareeshan — Optimizing Machine Learning Models with CUDA: A Comprehensive Performance Analysis — mixed precision, fusion, and library choice for ML.
• CUDA C++ Best Practices Guide — the canonical optimization checklist.
• NVIDIA — Occupancy API & launch configuration — using the calculator from code.