ESSAY 10 · 2026.04.19 · L4

Can 50 lines of Python replace 5000 lines of CUDA?

Where does the cost of abstraction show up? Four kernels on L4, four regimes with the answer.

The question

Triton's pitch is "CUDA-level performance without CUDA." Sounds like marketing — but Tri Dao actually wrote Flash Attention-2 in Triton. OpenAI's Kernel Gym is Triton. Is the claim true?

I decided to answer only with numbers. The four kernels from Lessons 1–6 — reduction, softmax, matmul, flash attention — rewritten in Triton. Run CUDA / torch built-in / Triton side by side on the same L4 (sm_89). One question: in which regime is abstraction expensive, and in which is it free?

Four regimes

Evidence 1 · HBM-bound is a tie

82–86% of HBM floor (300 GB/s). The spread among the three approaches is noise. In this regime, everything is "who can drink HBM fastest," and there's no reason a JIT would fall behind hand-written CUDA.

Evidence 2 · Triton edges out cuBLAS in compute-bound

2.9× over our WMMA. More importantly, narrowly over cuBLAS. The reason: autotune configs explore more points than a human. And TF32/FP16 selection happens automatically through tl.dot's dtype.

Evidence 3 · Flash Attention (N=8192)

100 lines of Triton at 79% of cuDNN FA-2. This is why Tri Dao wrote FA-2 in Triton. Look at the number and Triton stops being an experiment.

One line = dozens of lines, repeated four times

• tl.sum(x, axis=0) = 15 lines of warp shuffle boilerplate
• tl.dot(a, b) = 50 lines of WMMA fragment + load_matrix_sync + mma_sync (with automatic TC selection by dtype)
• Grouped program-id swizzle in 9 lines — writing the same thing in CUDA is painful
• Online softmax 50 lines (Lesson 6) → tl.max + tl.maximum + tl.exp + tl.sum, 15 lines

Where the cost is real

At n=2²⁰ (4 MB) in the reduction of Lesson 3, Triton was 3–12× slower than CUDA. One cause: launch floor 50–100 µs. Python interpreter → autotune cache lookup → JIT (once) → cuLaunchKernel. The floor's relative weight grows as the kernel shrinks.

If a Transformer layer is ≥1 ms, 100 µs is <10% — fine. But launch 30 element-wise ops individually in Triton and you're done. For small ops, PyTorch eager or torch.compile is better.

Two footguns

(1) TF32 benchmark lie

torch.matmul(fp32) doesn't use TF32 by default. tl.dot does. Compare them head-to-head and you see "Triton beats torch 2×." A fair comparison:

torch.backends.cuda.matmul.allow_tf32 = True
torch.backends.cudnn.allow_tf32 = True

With those two lines, both sides take the TF32 TC path. Before I set them, Phase 3's 2× lead was a lie.

(2) Autotune stale writes

@triton.autotune can leave a partial buffer polluted by other configs mid-trial. Reduction partials got mixed, rel_err blew up to 1.75. Fix:

@triton.autotune(configs=..., reset_to_zero=["partial_ptr"])
@triton.jit
def kernel(...): ...

best = kernel.best_config
block = best.kwargs["BLOCK_SIZE"]
return partial[:cdiv(n, block)].sum()

So why still CUDA

1. Triton hits walls — Blackwell's BF8/FP4 mma, persistent kernel, async-copy fine-grained control.
2. Reading PTX — chasing perf bugs means reading the PTX Triton emits. *.ptx, *.cubin in TRITON_CACHE_DIR.
3. Reference code — vLLM, FA-3, Mamba kernels are still CUDA. Reading them requires CUDA as your first language.
4. Perf diagnosis vocabulary — "why is it slow?" answers are bank conflict, register spill, occupancy. Those concepts carry into Triton unchanged.