Qwen Releases FlashQLA: 2-3x Faster Linear Attention Kernels on Hopper

Qwen has open-sourced FlashQLA, a high-performance linear attention kernel library built on TileLang that targets the GDN (Gated Delta Network) Chunked Prefill workload. Released in late April 2026 under the MIT License, FlashQLA reports a 2–3× forward-pass speedup and roughly 2× backward-pass speedup over the FLA Triton kernels, and roughly 2× over FlashInfer 0.6.9 on the same shapes — pushing linear-attention prefill closer to the bandwidth and Tensor Core limits of Hopper-class GPUs.

Advanced

What FlashQLA Targets

Linear attention variants such as Gated Delta Networks (GDN) replace the quadratic softmax attention with a recurrent state that, in principle, is much cheaper at long context. In practice the prefill phase — where the model ingests a long input prompt — is dominated by chunked matrix products with awkward shapes for tensor cores: small head dimensions, long sequences, and many gating operations sandwiched between GEMMs. Existing Triton implementations from Flash Linear Attention (FLA) and the general-purpose FlashInfer kernels leave a sizable performance gap on Hopper hardware.

FlashQLA is a focused rewrite of GDN Chunked Prefill — both forward and backward — that fuses the gate, decay, and matmul steps into warp-specialized kernels written in TileLang. The library requires SM90 or above (Hopper, H100/H200, and SM121 GB10 systems such as DGX Spark), CUDA 12.8+, and PyTorch 2.8+.

Three Optimizations Doing the Work

According to the project README, three design choices drive the speedup:

1. Intra-card context parallelism under tensor parallelism. By exploiting the exponential decay property of the GDN gate, FlashQLA automatically enables intra-card CP under TP, long-sequence, and small-head-count settings, improving GPU SM utilization in regimes where a single TP shard otherwise cannot saturate the device.
2. Algebraically reformulated forward and backward flows. The Chunked Prefill math is rewritten to reduce Tensor Core, CUDA Core, and SFU (special function unit) overhead while preserving numerical precision — important because GDN’s recurrent state can amplify rounding error if the kernel cuts corners.
3. Warpgroup-specialized fused kernels in TileLang. Rather than chaining many independent kernels, or fusing the entire computation into a single mega-kernel, FlashQLA hand-partitions a few key fused kernels with manual warpgroup specialization so that data movement, Tensor Core math, and CUDA Core math overlap cleanly.

Benchmarks

The repository’s benchmarks sweep head counts h_k,v ∈ {64, 48, 32, 24, 16, 8} — corresponding to TP1 through TP8 sharding for a 64-head model — against FLA 0.5.0 (Triton 3.5.1), FlashInfer 0.6.9, and TileLang 0.1.8 baselines. Across that sweep, FlashQLA reports a 2–3× forward-pass speedup over the FLA Triton kernels and around 2× on the backward pass. Smaller per-shard head counts, where SM utilization is hardest to maintain, are where FlashQLA’s intra-card CP path opens the largest gaps.

Why It Matters

FlashQLA is narrow on purpose: it is not a drop-in replacement for FlashAttention-style softmax kernels, and it is not a general inference engine. It is a kernel library for one specific shape of model — linear-attention LLMs that use GDN gating and chunked prefill, such as recent Qwen architectures — on one specific class of GPU. But that is exactly the gap where Triton-based reference kernels were leaving the most performance on the table, and it is the path Qwen has been pushing on for long-context efficiency. Pairing better kernels with the broader Qwen3.x stack means cheaper, faster prefill on long prompts without changing the model itself.

It is also another datapoint for TileLang’s trajectory. Where FlashAttention-4 leaned on hand-tuned CUTLASS for Blackwell, FlashQLA shows TileLang reaching production-quality results on Hopper for a non-trivial new operator class — with the warpgroup specialization expressed at a higher level than raw CUDA.

Related Coverage

• FlashAttention-4: Algorithm and Kernel Co-Design for Blackwell GPUs — the softmax-attention counterpart, targeting Blackwell instead of Hopper.
• Qwen3.6-27B: A Dense 27B Model That Beats a 397B MoE on Coding — the model family whose architecture FlashQLA is tuned to accelerate.
• Luce DFlash Brings 2x Speculative Decoding to Qwen3.6-27B on a Single RTX 3090 — a complementary inference-side speedup on the decode path.

This post was drafted with AI assistance and reviewed by RITS staff.

Sources

• QwenLM/FlashQLA — GitHub repository
• NVIDIA Developer Forums — Qwen introduces FlashQLA
• TileLang — tile-ai/tilelang