LESSON 08 · 2026.04.19 · L4
Triton vs CUDA — where does the cost of abstraction show up?
Four kernels from lessons 1–6 (reduction, softmax, matmul, flash attention) rewritten in 40–130 lines of Triton. Four regimes measured on L4 (sm_89).
Hardware upgrade — why L4
T4 has only one FP16 WMMA flavor and no TF32. To measure Triton tl.dot automatically picking TF32, you need a GPU with TF32 Tensor Cores. L4 (Ada) = TF32 121 · FP16 242 · FP8 485 TFLOPS, L2 48 MB (8× T4's).
Memory-bound — three approaches tie within 10%
| task | CUDA | torch | Triton |
|---|---|---|---|
| Reduction 67M fp32 | 258 GB/s | 254 GB/s | 245 GB/s |
| Softmax 4096² fp32 | 237 GB/s | 240 GB/s | 221 GB/s |
Triton lands at 93–95% of hand-written CUDA. When HBM is the bottleneck, the abstraction tax evaporates.
Compute-bound — Triton narrowly beats cuBLAS
| task | CUDA | torch | Triton |
|---|---|---|---|
| matmul 4096³ FP32 (TF32) | 3.9 TF | 25.8 | 28.9 |
| matmul 4096³ FP16 | 18.5 TF | 51.8 | 54.0 |
2.9× over our WMMA. Over cuBLAS: FP32 +12%, FP16 +4%. The spot where autotune narrowly edges human hand-tuning.
Flash Attention — Triton fp16 is polyglot
| N | CUDA FA fp32 | Triton fp32 | Triton fp16 | SDPA fp16 |
|---|---|---|---|---|
| 1024 | 0.324 | 0.148 | 0.122 | 0.076 |
| 2048 | 0.638 | 0.196 | 0.138 | 0.076 |
| 4096 | 1.256 | 0.358 | 0.207 | 0.127 |
| 8192 | 3.045 | 1.118 | 0.496 | 0.394 |
N=8192: Triton fp16 is 6.14× over our CUDA FP32, and 79% of SDPA (cuDNN FA-2). 100 lines of Triton reaches 80% of cuDNN. This is why Tri Dao wrote FA-2 in Triton.
One line = dozens of lines, repeated four times
tl.sum(x)= 15 lines of warp shuffle boilerplate from Lesson 3tl.dot(a,b)= 50 lines of WMMA fragment + mma_sync from Lesson 5 (plus automatic TC selection by dtype)- Grouped program-id swizzle — 9 lines here, painful to write in CUDA
- Online softmax 50 lines (Lesson 6) →
tl.max + tl.maximum + tl.exp + tl.sum, 15 lines
Regions where the abstraction cost is real
| region | Triton vs CUDA | cause |
|---|---|---|
| Small N (< 4 MB) | 3–12× slower | Launch floor 50–100 µs (Python → autotune cache → JIT → cuLaunch) |
| HBM-bound | 95% | almost none |
| Large matmul / FA | wins | autotune picks a better config than a human |
Practical line: if the Transformer layer is ≥1 ms, a 100 µs overhead is <10% — tolerable. If you launch 30 element-wise ops individually in Triton, you're cooked.
Two footguns
(1) TF32 benchmark lie. torch.matmul(fp32) doesn't use TF32 by default. tl.dot does. Compare them head-to-head and "Triton beats torch 2×" — which looks great but is wrong. For a fair fight:
torch.backends.cuda.matmul.allow_tf32 = True
torch.backends.cudnn.allow_tf32 = True(2) Autotune stale writes. @triton.autotune trials can leave tainted writes in the output buffer. Use reset_to_zero=["partial_ptr"] and slice with best_config.kwargs["BLOCK_SIZE"] after the call.
Why keep learning CUDA
- Triton hits a wall and you fall back to CUDA (Blackwell mma, persistent kernels, async copy).
- You need to read the PTX Triton emits to chase perf bugs.
- vLLM, FA-3, Mamba are still CUDA.
- The answer to "why is this slow" is bank conflict, register spill, occupancy — all CUDA concepts.
CUDA = assembly, Triton = C. Write most of it in C; reserve assembly for hot paths.