Requirements

Real reason is that the speedup comes from a producer/consumer warp split that uses TMA (Tensor Memory Accelerator) plus WGMMA (warpgroup MMA) plus async barriers. SM90 (H100) was the first arch to ship all three, and SM100 (Blackwell) extends them.

Quick map:

• TMA does async bulk loads from HBM straight into shared memory with one PTX instruction, so producer warps fire-and-forget tile loads while consumer warps run matmul on already-loaded tiles. SM89 (4090 / Ada) has no TMA, so async loads fall back to cp.async with 16-byte granularity and per-thread bookkeeping, which kills the latency hiding budget for a 16-stage pipeline.
• WGMMA is a warpgroup-scoped tensor core instruction that operates on 4 warps at once and reads operands directly from shared memory. Pre-Hopper tensor cores (mma.sync) need operands in registers, so you spend register file on staging. With WGMMA you free those registers for accumulator state, which is what lets a deep pipeline fit in 228KB of SM SRAM.
• The 16-stage warp-specialized pipeline line is the FlashAttention-3 (Shah 2024) playbook applied to gated-DeltaNet backward. The backward pass reverse-scans through the recurrent state, you need both forward states and gradients live simultaneously, and the only way to keep that in SRAM at long context is the deep async pipeline that WGMMA + TMA permits.

Could it be ported to SM89? Yes, but the speedup mostly evaporates. You'd be back to FlashAttention-2 style cooperative load with mma.sync, register pressure goes up, pipeline depth has to come down. At that point Songlin Yang's FLA Triton kernels (flash-linear-attention library, covers GDN, GLA, RWKV-7, Mamba2) are already a reasonable baseline on Ampere/Ada, within roughly 30% of theoretical max for those tiers. The 2-3x in the post is a headline against FLA on SM90, not a hardware-agnostic number.

Practical: matters for Qwen3-Next style hybrid models (mostly linear-attention layers with a few full-attention layers interleaved), at >32k context, on H100/H200/B200. On a vanilla full-softmax Qwen3 it does nothing; on a 4090 it's not what you want.