Triton — the pragmatic kernel language for inference¶

Version target (July 2026): Triton 3.7.1. Follow: https://github.com/triton-lang/triton/releases

Triton is where you will spend the bulk of your kernel-writing hours in Phases 3–4. It gives you ~90% of hand-rolled CUDA performance for ~10% of the effort, integrates natively into PyTorch (via torch.compile), and is the language of choice for vLLM, SGLang, and every serious inference stack when they need a custom kernel fast.

The Triton mental model¶

You write Python-decorated kernels. The compiler lowers them to LLVM IR → PTX/AMDGPU.
Threads are hidden. You program at the block level. tl.program_id, tl.arange, tl.load operate on blocks/tiles, not scalar threads.
You still own tile sizes and memory hierarchy. BLOCK_M, BLOCK_N, BLOCK_K, num_stages, num_warps are yours. Pick well.
Autotuning is a first-class feature. @triton.autotune sweeps configs, caches best.
The compiler handles SMEM allocation, coalescing, and pipelining — mostly. When it doesn’t, you’ll drop into inspecting the generated IR (ttir, ttgir, ptx).

What changed since 2024¶

Triton has moved fast:

Triton 3.6+ (2025): first-class Blackwell (sm100) support — tcgen05-lowering paths, TMEM utilization for MMAs.
Triton 3.7 (2026): AMD RDNA4 (gfx1250) support, warp specialization on AMD, Tensor Data Movement (TDM), warp-pipeline path.
Gluon (new, ~9 months old as of July 2026): OpenAI’s lower-level DSL sharing the Triton compiler stack but exposing explicit tensor layouts. Think “Triton without the auto-layout inference.” See https://www.lei.chat/posts/gluon-explicit-performance. Use case: when Triton’s auto-layout chooses badly and you need to force a specific SMEM/register layout.
The is_hopper / is_blackwell capability checks are standard in modern tutorials; the fused-attention tutorial has hardware-specific paths.

Tutorial order — do these in this exact sequence¶

Official tutorials live at https://triton-lang.org/main/getting-started/tutorials/. Do them in order, don’t skip:

#	Tutorial	What you learn	Rough time
01	Vector Addition	`tl.program_id`, `tl.arange`, block-level thinking	30 min
02	Fused Softmax	Reduction, masking, `tl.max`/`tl.sum` on tiles	1–2 h
03	Matmul	Autotune, `tl.dot`, pipelining stages, `num_stages`/`num_warps`	3–4 h
04	Low-memory Dropout	RNG in kernels, in-place ops	30 min
05	LayerNorm / RMSNorm	Multi-pass reduction, welford, gradient stubs	2 h
06	Fused Attention	THE tutorial — online softmax + FA2 in Triton, Hopper/Blackwell branches	4–6 h
07	Grouped GEMM	Batched-with-varying-shapes matmul	1 h

Tutorial 06 is the single most important artifact you’ll produce in Phase 2/3. Everything in Phase 3 rides on your ability to read and modify it. Do it twice: once following, once from scratch.

Idiomatic patterns you must own¶

Autotune template¶

import triton
import triton.language as tl

@triton.autotune(
    configs=[
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 64,  'BLOCK_K': 32}, num_stages=3, num_warps=4),
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 128, 'BLOCK_K': 32}, num_stages=4, num_warps=8),
        triton.Config({'BLOCK_M': 64,  'BLOCK_N': 64,  'BLOCK_K': 64}, num_stages=3, num_warps=4),
    ],
    key=['M', 'N', 'K'],
)
@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K,
                  stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn,
                  BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
    ...

Key moves:

key=[...] triggers a re-tune only when those args change.
num_stages = software pipelining depth. Higher = more SMEM use, more overlap.
num_warps = warps per block. 4–8 typical.

Masking template¶

mask_m = offs_m < M
mask_n = offs_n < N
a = tl.load(a_ptr + offs, mask=mask_m[:, None] & mask_k[None, :], other=0.0)

The other=0.0 value matters for numerics — for softmax you’ll use -inf.

Reduction across a dim¶

row_max = tl.max(x, axis=1)     # per-row max
row_sum = tl.sum(tl.exp(x - row_max[:, None]), axis=1)
y = tl.exp(x - row_max[:, None]) / row_sum[:, None]

This is the fused softmax core. You will re-derive this in ../04_attention/01_online_softmax.md.

Common pitfalls (real ones, from real bugs)¶

tl.constexpr vs runtime arg. Block sizes must be constexpr (compile-time). If you accidentally pass them as runtime ints, you get wrong specialization + terrible performance.
other= value on masked loads. Wrong other value → wrong reductions. For max, use -inf; for sum, use 0.
Autotune cache poisoning. If your key=[...] doesn’t include a dimension that matters (e.g., dtype changed), you’ll get a bad config on the wrong shape. Include everything that affects layout.
num_warps too high. More warps ≠ faster. Pick 4 or 8 unless you have reason otherwise; larger num_warps starves per-thread registers.
num_stages too high. Stages consume SMEM; on smaller GPUs you’ll spill or fail to allocate silently.
Forgetting tl.static_assert. Assert BLOCK_K % 16 == 0 etc. so misconfigurations fail loudly.
Debugging without TRITON_INTERPRET=1. Triton’s interpret mode runs the kernel on the CPU as pure Python — excellent for debugging shape/mask logic.
Measuring with time.time(). Use triton.testing.do_bench or torch.cuda.Event. time.time() includes launch overhead + JIT compilation.
Ignoring TRITON_PRINT_AUTOTUNING=1. Turn this on when developing — you’ll see which config won.
Not inspecting IR. For hard performance work: TRITON_CACHE_DIR=... TRITON_ALWAYS_COMPILE=1 then read .ttgir / .ptx from the cache. This is where you find out whether your loads actually vectorized.

When to reach for Gluon (not yet, but know it exists)¶

Gluon is lower-level Triton — you specify tensor layouts explicitly (blocked, mma, wgmma, tcgen05 layouts). Use case: when Triton’s autoinference picks a bad layout and you can’t nudge it via block sizes / num_warps. This is post-Phase 4 material — don’t touch it before you’ve hit the ceiling of vanilla Triton.

Docs: https://triton-lang.org (Gluon section). Explainer: https://www.lei.chat/posts/gluon-explicit-performance

Modern examples to read (all Triton, all inference)¶

vLLM Triton kernels: https://github.com/vllm-project/vllm/tree/main/vllm/_custom_ops — PagedAttention, RoPE, RMSNorm, quant
Unsloth Triton: https://github.com/unslothai/unsloth — fused fine-tuning kernels
Liger-Kernel (LinkedIn): https://github.com/linkedin/Liger-Kernel — fused SwiGLU/RMSNorm/RoPE/cross-entropy for training but reusable
FlagAttention (BAAI): reference implementations of attention variants
Triton fused attention tutorial itself is a production-grade example

Benchmark hygiene (repeat with me)¶

Warm up the kernel (3–5 calls before measuring).
Use torch.cuda.synchronize() before/after the timed region or use CUDA events.
Report median of ≥ 100 iterations, plus p95/p99 for latency-sensitive work.
Fix clocks if reproducibility matters: nvidia-smi -lgc <freq>.
Compare against both torch.nn.functional.<op> and the hand-rolled cuBLAS/cuDNN if applicable.
Verify correctness first, benchmark second. torch.allclose(y, ref, atol=1e-2, rtol=1e-2) for bf16.

What “done with Triton” looks like for Phase 2¶

You’ve completed tutorials 01–06.
You’ve written a fused softmax + a fused RMSNorm from scratch (no peeking at the tutorial), both faster than eager PyTorch by ≥ 2× on your GPU.
You’ve run triton.testing.do_bench correctly on all of them.
You know how to inspect the generated PTX to check whether a load vectorized.

That unlocks Phase 3, where you’ll write FA2 in Triton (../04_attention/04_writing_fa_in_triton.md).

References¶

Official docs + tutorials: https://triton-lang.org/main/getting-started/tutorials/
Fused Attention tutorial: https://triton-lang.org/main/getting-started/tutorials/06-fused-attention.html
Releases + changelog: https://github.com/triton-lang/triton/releases
OpenAI Triton launch post (still worth reading): https://openai.com/index/triton
Gluon overview: https://www.lei.chat/posts/gluon-explicit-performance
vLLM custom ops: https://github.com/vllm-project/vllm/tree/main/vllm/_custom_ops
Liger-Kernel: https://github.com/linkedin/Liger-Kernel