Last night’s Dist-MoE note was Meta’s Blackwell SM100 MoE kernels — fused dispatch, grouped expert GEMM, and combine with a named BF16 / MXFP8 / NVFP4 matrix. Yesterday AM’s Kolibri-1 was open serving MoE weights. Olmo-core 3 answered the open MoE trainer. Tonight is the complementary attention piece: on October 1, 2026, the PyTorch blog published Optimizing Jagged Flash Attention with TLX — Meta’s production jagged / broadcast-Q attention kernel for the Generative Ads Model (GEM), rewritten in TLX (Triton Low-level Extensions) for NVIDIA B200, with a named comparison against open-source FlashAttention-4.
Primary:
PyTorch blog
·
GitHub tlx_jfa.
Every number below is Meta’s, from the blog and linked repo, measured on their shapes. Meta did not publish tok/s or MFU, so I don’t either.
Spec sheet (named numbers)
| TLX Jagged Flash Attention (JFA) | |
|---|---|
| Published | 2026-10-01 (PyTorch blog; “Last updated: 2026-10-01”) |
| What it is | Jagged Flash Attention for Meta GEM on NVIDIA Blackwell B200, authored in TLX |
| Workload | Packed variable-length (jagged) sequences + offsets; production case broadcast-Q (broadcast_q=True) — one dense Q against jagged K/V |
| Baseline | Open-source FlashAttention-4 (May 2026 CuteDSL) |
| Code size (Meta claim) | TLX JFA ~3.2K lines Triton-level vs FA4 ~10K lines CuteDSL (~3× less) |
| Jagged comparison (bf16, B200) | Forward ~+13% avg vs FA4; backward ~+50% avg vs FA4 |
| Dense comparison (bf16, B200) | Forward ~87% of FA4; backward ~+17% vs FA4 |
| Dense shape cited | B=768, H=4, head_dim=128 |
| Precision in main comparison | bfloat16 |
| Structural | Warp specialization; explicit SMEM/TMEM; producer/consumer barriers; persistent CTAs; CLC (Cluster Launch Control) |
| Soft balancing (fwd) | Host sort + zigzag tile→SM schedule — Meta claims ~20% forward recovery vs naive mapping |
| dQ epilogue | Multi-stage double-buffered SMEM staging; early TMEM release recovered +8–11% tensor-core utilization in waterfall tests |
| 2-CTA MMA (bwd) | Paired-CTA tcgen05 on broadcast-Q HEAD_DIM=128 — Meta claims ~+12% throughput / −11% latency vs single-CTA |
| Forks documented | MXFP8 (E4M3 + E8M0 block scales); block-sparse top-k (~1.3–1.5× forward vs dense at 0.5 selection) |
| Code |
tlx_jfa
under ads_model_kernel_library
|
The question this post answers is narrow and useful: can you beat FA4 on the jagged broadcast-Q shapes that dominate a production ads ranking model, in readable TLX rather than 10K lines of CuteDSL? Meta says yes, and shows the work.
What it is selling
Dist-MoE sits under routed experts once you already have IDs and scores. TLX JFA sits under attention on packed jagged sequences — the kernel Meta says is the single slowest piece inside GEM. Padding those sequences to a fixed length can waste up to 50% of compute (Meta’s earlier GEM writeup, cited from the Oct 1 post). JFA runs FlashAttention directly on contiguous packed Q/K/V plus an offsets tensor, without materializing padded tokens.
Plain Triton left on-chip movement and scheduling to the compiler. TLX exposes the knobs Meta needs on Blackwell: explicit SMEM/TMEM allocation, async_task warp specialization, barriers, async TMA / MMA, and Cluster Launch Control. The CTA splits into role-specialized warps — TMA loads, tensor-core matmuls, softmax/correction, epilogue store, plus a dedicated dQ-reduction warp in the backward — so tensor cores stay on back-to-back MMAs while softmax runs concurrently.
The production case shapes the design: a single dense Q is broadcast across every sequence in the batch, so the backward dQ must be summed across the entire batch. That turns the dQ epilogue into a contended cross-program reduction — the #1 backward bottleneck Meta’s Triton-MPP barrier analysis pinned at roughly 9–11% lost tensor-core utilization before the multi-stage staging fix.
Structural changes and layered opts (Meta’s list)
Structural: warp specialization + explicit SMEM/TMEM + barrier pipelining; persistent CTAs (one per SM); triple-buffered K/V; TMEM aliasing across non-overlapping lifetimes.
On top of that structure:
-
Jagged tile scheduling. Forward: sort tiles by KV workload, zigzag across SMs (~20% forward recovery claimed). Backward: precompute valid
(batch, head, K/V-block)tiles, round-robin for L2 reuse, pad to cluster size for the 2-CTA path. Residual runtime variance handed to CLC. - Multi-stage dQ staging. Double-buffered SMEM pipeline so one slice’s HBM reduce-add overlaps the next slice’s TMEM→SMEM copy.
- Early TMEM release. Pre-load the last 1–2 dQ slices into registers, free TMEM, let the MMA warp start the next tile’s dQ matmul while stores finish — waterfall tests recovered +8–11% TC utilization.
- Loop peeling. Branch-free bulk KV pass + tiny masked tail; Meta’s point is register pressure, not the cheap branch itself.
- 2-CTA collaborative MMA (backward). Two CTAs cooperate on a wider matmul over adjacent K/V blocks; Meta claims ~+12% throughput / −11% latency on broadcast-Q HEAD_DIM=128.
Where it sits next to earlier posts
| Layer | Project | What it gives you |
|---|---|---|
| Blackwell MoE kernels (last night) | Meta Dist-MoE | SM100+ fused dispatch + grouped GEMM + combine; BF16/MXFP8 train+infer; NVFP4 infer |
| Open serving MoE weights (yesterday AM) | Kolibri-1 | 78.1B/3.46B, hybrid SWA, ~78 GB FP8 + vLLM plugin |
| Open MoE trainer | Ai2 Olmo-core 3 | B300 capacity tax fix; training tok/s |
| Zero-full-attn SWA/DSA decode | Naive-N0.5-Flash | Native 1M with no full-attention layers |
| Near-package LPDDR fanout | Credo Weaver | Claimed HBM4-class edge BW at ~25× capacity |
| Blackwell attention kernels (this post) | Meta TLX JFA | B200 jagged / broadcast-Q FA; ~+13% fwd / ~+50% bwd vs FA4; ~3.2K TLX lines |
This isn’t a Dist-MoE follow-up. Different layer: attention vs MoE. Dist-MoE assumes you already routed experts. TLX JFA assumes you already packed jagged sequences. Neither ships a checkpoint.
Related but separate: Meta’s earlier TLX Block Attention post (2026-07-21) targets fixed block-diagonal attention with different benches (FA2 baselines, fused rotary). The October 1 post is the jagged / broadcast-Q JFA vs FA4 comparison, and the two shouldn’t be confused.
Caveats
- Jagged production shapes, not a universal FA4 killer. Headline wins are on jagged broadcast-Q. On dense equal-length shapes, forward is ~87% of FA4 while backward still wins (~+17%). The jagged +50% backward number is not a dense LLM result.
- Vendor / shape-owned benches. All comparison percentages are Meta’s bf16 B200 measurements against May 2026 open-source FA4 on their shapes and sparsities (0.1 → 0.9 jagged sweep). Re-run on your shapes before you plan capacity.
- Blackwell / TLX stack. Targets B200 with TLX primitives (async TMA/MMA, TMEM, CLC). Ampere / Hopper and plain-Triton-only stacks are outside the published comparison.
- Not Dist-MoE. Dist-MoE is CuTe DSL MoE kernels (ARM64 + SM100, named precision matrix, no published throughput benches). TLX JFA is Triton/TLX attention with published FA4 comparison percentages. Different API, different layer.
- Ads GEM context. Motivating workload is Meta’s Generative Ads Model (and Kunlun, cited). Uniform dense LLM decode may not share the jagged / broadcast-Q bottleneck.
- MXFP8 / sparse forks are not the main comparison table. Main FA4 comparison is bf16. MXFP8 and block-sparse results are separate forks Meta describes qualitatively — treat as Meta claims on their forks.
How I would evaluate
- Read the Oct 1 PyTorch post — especially the jagged vs dense charts and the broadcast-Q dQ discussion — before quoting any percentage.
-
Clone
tlx_jfaand pin the TLX / Triton / CUDA stack the repo documents; do not assume Dist-MoE’s ARM64+cu130 pins apply. - Reproduce the jagged broadcast-Q sweep (sparsity 0.1 / 0.5 / 0.9) and the dense B=768 H=4 head_dim=128 case against the same May 2026 FA4 reference; log NCU tensor-core / TMEM counters the way the post describes.
- Exercise MXFP8 and block-sparse forks only after the bf16 jagged path is clean.
- Keep this separate from Dist-MoE’s MoE precision matrix, Kolibri’s serve residency, Olmo-core’s trainer tok/s, and Naive-N0.5-Flash’s zero-full-attn decode path — this is Blackwell jagged attention.
Bottom line
If you run Blackwell training stacks and your attention traffic is jagged or broadcast-Q, Meta published the comparison numbers: TLX JFA claims ~+13% forward and ~+50% backward vs May 2026 FA4 on those shapes, in ~3.2K lines of readable TLX instead of ~10K CuteDSL, with open code under tlx_jfa. Last night answered MoE kernels. Tonight answers attention. Same generation, different layer — and the layer most GEM-like ranking models actually stall on.
Create-only. Specs from Meta’s October 1, 2026 PyTorch blog on Optimizing Jagged Flash Attention with TLX and the facebookresearch/ads_model_kernel_library tlx_jfa tree. Update when third-party rematches, tok/s, or MFU numbers ship.
Sources:
PyTorch blog
·
GitHub tlx_jfa.
Anpoo Sivanadi, Staff Software Engineer