A Triton attention forward kernel within 5% of FlashAttention-3 on H100
Difficulty 20, the lowest in the pack: a Triton-based compiler (Tawa) already reports up to 96% of FlashAttention-3 in FP16 on a development branch, so what remains is BF16 causal coverage and landing the feature in a release. Application value 49: parity would let researchers write new attention variants in Python at full speed. Leverage 35: higher than the GEMM parity tasks because it tests whether a compiler can express Hopper warp specialisation, which every later Triton kernel inherits.
The problem
Can attention written in Triton reach 95% of FlashAttention-3 on H100? A Triton-based research compiler is close.
- 1FlashAttention-3 is hand-written low-level C++
- 2Triton is Python-level, usually slower
- 3Compiler must split warps into roles
- 4Tawa branch reaches 96% in FP16
- 5Released Triton at 95%, BF16 causal
What is asked. Attention is the step in a transformer where each token compares itself with earlier tokens and takes a weighted average of their information. FlashAttention computes it tile by tile without ever storing the huge table of scores (see the FlashAttention entry). Its third version, FlashAttention-3, is written in low-level C++ for NVIDIA's H100 and reaches up to 740 TFLOP/s in FP16. Triton is a Python-level language in which a kernel is written for whole tiles and a compiler handles the details. The problem: a Triton attention forward kernel, using only a released version of Triton, at 95% or more of FlashAttention-3's speed for one fixed case. The case is BF16 numbers, head dimension 128, causal masking (a token may only look backwards), and sequences of 8192 tokens.
Path to understanding
foundations → problemThe research entries this problem rests on, in the order to read them. Each links to its explainer.
- Start with
- Step 2
- Step 3
- FlashAttention: IO-aware exact attentionAlgorithm · host
- Then
- A Triton attention forward kernel within 5% of FlashAttention-3 on H100The problem
State of the art
FlashAttention-3 is written in CUTLASS-style C++ with TMA, WGMMA, warp specialisation and ping-pong scheduling, and its authors report up to 740 TFLOP/s in FP16 on H100 (arXiv abstract). Triton includes a fused-attention tutorial kernel, and releases since 3.2 add TMA descriptors and automatic warp specialisation on Hopper; PyTorch's warp-specialisation post reports gains of 10–15% for FlashAttention-style kernels from it, not parity. The closest result is Tawa (Chen et al., arXiv:2510.14719, CGO 2026), an automatic warp-specialisation compiler built on Triton, which reports up to 96% of FlashAttention-3 throughput for FP16 attention forward on H100 (89% in FP8), from a development branch rather than a Triton release, and without a BF16 causal result at sequence 8192. Triton's lower-level Gluon dialect ships an attention example only for Blackwell. As of 2026-10-17 no released Triton kernel is documented at 95% of FlashAttention-3 on the configuration below.
Realistic applications
A compiler-language kernel at parity would make attention variants (new masks, positional schemes, sparsity patterns) cheap to write without sacrificing speed, and would indicate that the compiler captures Hopper asynchrony.
What counts as solved
A public Triton kernel at a pinned Triton release that meets the 95% bar on all configurations in the spec, reproduced by a second party.
How it gets verified
GPU benchmarkA kernel run and timed on stated hardware against a baseline.
Hardware: one H100 SXM 80 GB. Shapes: head dim 128, causal, BF16, sequence 8192 with batch x heads chosen so total tokens stay near 16384 (for example batch 2, 32 heads), plus sequence lengths 2048 and 16384 as secondary reports. Baseline: FlashAttention-3 from the Dao-AILab/flash-attention hopper/ directory at a pinned commit, compiled with CUDA 12.8 or later. Metric: forward TFLOP/s using the causal FLOP convention of the FlashAttention benchmarks, median of at least 100 iterations after warmup. Correctness: max absolute error versus an FP32 PyTorch reference no larger than 2x the error of FlashAttention-3 on the same inputs. Harness: a public script at a pinned commit timing both kernels in one process with triton.testing.do_bench or CUDA events.
Claims under verification
+ add a claimNo claimed results on the Verification board for this problem yet.
Where the work is
Compiler and scheduling engineering; some new ideas needed to expose asynchrony in a tile-level language.
Bottlenecks
ranked by votes- #1EngineeringExpressing warp specialisation in the language
Hopper-level performance relies on explicit producer and consumer warps, register reallocation and pipeline barriers that a tile-level language must infer or expose.
Community grades
Sign in to add your own grade.
Connected problems
- This problem is related to An open attention backward kernel within 10% of FlashAttention-3 for head dimension 256Same hardware and baseline; the backward pass is harder because of atomics and recomputation.
Progress
- Oct 16, 2025Tawa reaches up to 96% of FlashAttention-3 (FP16)source
A Triton-based compiler with automatic warp specialisation reports near-parity with FlashAttention-3 for FP16 attention forward on H100, from a development branch.