MpFA: Hardware-Efficient Train-Free QK4V8 FlashAttention Kernels on Blackwell GPUs Chencheng Deng, Jianbin Fang, Dezun Dong* College of Computer Science and Technology, National University of Defense Technology Changsha, China {chenchengdeng,j.fang,dong}@nudt.edu.cn Attention
75 50 73 57
25 26
0
Others
100
40
8
11
1K
4K 16K 32K 64K 128K
context length
(a) prefill
GPU time share (%)
Long-context LLM inference pushes modern GPU serving stacks into an attention-bound regime, where both compute and memory are dominated by the softmax–GEMM pipeline. On NVIDIA Blackwell GPUs, FP4 Tensor Cores offer high matmul throughput, but we find that fully FP4 attention often fails to translate this throughput into end-to-end speedups due to non-matmul costs: online quantization after softmax, tensor/shared-memory data movement, and contention on the softmax path. We present MpFA, a training-free FlashAttention kernel optimized for Blackwell. Guided by hardware characterization, MpFA uses mixed precision: NVFP4 for 𝑄𝐾 ⊤ and FP8 for 𝑃𝑉 (QK4PV8). This preserves low-bit 𝑄𝐾 ⊤ throughput while avoiding the conversion and scaling overheads of FP4 𝑃𝑉 . To recover accuracy without further stressing the softmax pipeline, MpFA introduces rank-one smoothing compensation implemented as an additional Tensor Core MMA. MpFA further improves performance with a fine-grained asynchronous pipeline, tensor-memory reuse, and adaptive parallel partitioning across prefill and decode. On an NVIDIA B200 and across 16K–128K contexts, MpFA improves prefill throughput over state-of-the-art BF16/FP8 baselines and increases end-to-end output throughput by 2.81× over BF16 FA4 across Llama-3.1-8B and Qwen3-14B. Across five benchmark suites and two models, rank-one compensation recovers 62.5% of the accuracy loss with about 2.0% kernel overhead.
75 50 62
25
75
85
91
34 15
0 1K
4K 16K 32K 64K 128K
context length
(b) decode
Figure 1. Device-side GPU time breakdown for Llama-3.18B served by SGLang on a B200 with the BF16 FA4 backend.
the growing KV cache, resulting in memory traffic that increases linearly with context length and low arithmetic intensity. We run Llama-3.1-8B with SGLang’s standard BF16 FA4 backend on a B200. As shown in Figure 1, attention accounts for 72.5% of device-side GPU time in prefill and 91.4% in decode at 128K tokens. Although FlashAttention-4 (FA4) [27] pipelines computation through tensor memory and fifth-generation Tensor Cores, the evaluated backend uses BF16 and does not exploit FP4 Tensor Cores for the two dominant matrix multiplications. Quantization-aware training can adapt a model to 4-bit attention, but requires costly per-model retraining and is difficult to deploy at scale. Post-training methods avoid retraining; for example, SageAttention3 [30] applies NVFP4 attention to the consumer Blackwell RTX 5090, while HiFA4 [6] targets the Ascend HIF4 NPU. However, neither is designed to exploit the B200’s tensor memory, fifth-generation Tensor Cores, and operand-delivery paths or to jointly support decode optimization, paged KV caches, and CUDA Graph capture. Efficient training-free FP4 attention therefore remains unavailable on data-center Blackwell GPUs. Our hardware characterization and profiling reveal two challenges in building training-free FP4 FlashAttention for data-center Blackwell GPUs. First, when fifth-generation Tensor Cores execute block-scaled FP4 matrix multiply-add (MMA), matrix operands must be supplied from shared memory, whereas scale factors must come from tensor memory. Quantizing the softmax output 𝑃 to FP4 adds per-block amax computation, scale-factor generation, and data movement
Keywords: FlashAttention, Quantization, Blackwell GPU
1
GEMM
100
GPU time share (%)
arXiv:2609.33135v1 [cs.DC] 27 Sep 2026
Abstract
Introduction
Applications such as retrieval-augmented generation, code understanding, and multi-turn agents increasingly rely on long contexts, making long-context inference an important LLM-serving workload [1, 12]. Meanwhile, AI infrastructure is adopting the Blackwell architecture. NVIDIA’s data-center B200 combines fifth-generation Tensor Cores, native low-bit computation, and tensor memory for storing operands and intermediate results. LLM inference consists of prefill and decode. Prefill processes the entire input sequence, so its attention computation grows approximately quadratically with context length. Decode generates one token at a time and repeatedly accesses * Corresponding author.
1
between tensor and shared memory. These operations load the softmax path, preventing low-bit matrix-multiplication speedups from translating into end-to-end performance. Second, existing smoothing compensation methods place the compensation term on the softmax path. This extra elementwise work further loads the already saturated path, where it competes with other operations and offsets the benefit of FP4 computation. We present MpFA, a high-performance, training-free lowbit attention kernel for LLM inference on data-center Blackwell GPUs. We characterize three precision assignments: NVFP4 for both matrix multiplications; NVFP4 for 𝑄𝐾 ⊤ with BF16 for 𝑃𝑉 ; and NVFP4 for 𝑄𝐾 ⊤ with FP8 for 𝑃𝑉 . Our evaluation shows that QK4PV8 provides the best balance between arithmetic throughput and non-matmul overhead on the B200, making it the hardware-efficient choice for FA4. FP8 𝑃𝑉 retains the benefits of low-precision Tensor Core execution while avoiding the additional precision-conversion requirements of an FP4 probability matrix. BF16 𝑃𝑉 avoids these conversion requirements as well, but leaves Tensor Core throughput on the table because FP8 Tensor Cores provide twice the throughput of BF16. MpFA therefore adopts QK4PV8 mixed precision and optimizes the FA4 pipeline for both inference stages: prefill uses a fine-grained asynchronous pipeline and tensor-memory column reuse to co-locate scale factors with double-buffered accumulators; decode restores parallelism through GQA packing, KV partitioning, and column-split softmax while avoiding exponential evaluations on padded rows. To reduce smoothing overhead, MpFA introduces a rankone Tensor Core MMA compensation method that moves the correction from the saturated softmax path to otherwise underutilized Tensor Cores. It stores 𝐾 in NVFP4 and 𝑉 in FP8, consuming the low-bit KV cache without dequantization and reducing KV bytes per token to 1/2.56 of BF16. Experiments on the B200 show that, across 16K–128K, MpFA achieves average prefill speedups of 1.31–1.33× over BF16 FA4 and 1.15–1.17× over FP8 FlashInfer. Across 8K– 128K, its decode kernel achieves a 2.08× average speedup over split-KV-enabled BF16 FA4. Across ten model–task pairs spanning five standard benchmarks, rank-one smoothing recovers 62.5% of the accuracy loss on average, with about 2.0% average kernel overhead. MpFA achieves 2.81× the endto-end output throughput of BF16 FA4 across both models and 16K–128K, and by 1.06× over FP8 FlashInfer across both models and 32K–128K contexts. This paper makes the following contributions: 1. We systematically analyze how attention kernels on data-center Blackwell GPUs become limited by nonmatmul units, providing insight for optimizing other fused operators (§3). 2. We present MpFA, a high-performance QK4PV8 FlashAttention kernel for LLM inference, with tensor-memory
reuse and stage-specific optimizations for prefill and decode (§5). 3. We introduce a rank-one Tensor Core smoothing algorithm that moves compensation from the saturated softmax path to MMA, recovering accuracy while maintaining hardware utilization (§4).
2
Background
2.1
Blackwell GPUs
Blackwell is NVIDIA’s GPU architecture for next-generation AI computing. This paper focuses on the data-center B200 rather than the consumer RTX 5090. Both provide Blackwell Tensor Cores and low-bit formats, but differ in compute capabilities and operand delivery. RTX 5090 attention kernels are optimized for the consumer GPU’s matrix multiply-add (MMA) path, whereas the B200 introduces on-chip tensor memory and tcgen05.mma. Consequently, RTX 5090 implementations cannot fully exploit the B200, whose adoption requires attention kernels redesigned around its data path [13]. From hardware to kernels. The B200’s Tensor Cores, tensor memory, and shared memory jointly determine the onchip data path for attention. Figure 2 summarizes the execution path considered in this paper. In Figure 2(a), the asynchronous copy engine (TMA) first moves tiled 𝑄, 𝐾, and 𝑉 from global memory to shared memory. Tensor Cores then use tcgen05.mma instructions to compute 𝑄𝐾 ⊤ and 𝑃𝑉 and write the intermediate results to tensor memory. CUDA cores and special function units execute online softmax and directly read from and write to the score, probability, and output accumulators in tensor memory. The final results are written back to global memory. The SS and TS labels in the figure denote two operand-delivery paths; as discussed in §3.1, these paths constrain the precision assignment. Unlike traditional execution models, which keep accumulators in registers, this design uses tensor memory as separate onchip storage for results, allowing matrix multiplication and softmax to form an asynchronous pipeline with explicit role separation [27]. Figure 2(b) shows the capacity constraint: in our kernel configuration, each cooperative thread array (CTA) is allocated 128 rows × 512 columns, or 256 KB. Two attention score accumulators and two output accumulators, each with block size 128 and FP32 elements, consume all 512 columns. Block-scaled FP4 matrix multiply-add. FP4 cannot represent the full dynamic range of activations by itself, so the B200 uses block-scaled FP4 MMA. NVFP4 associates every consecutive 𝐺=16 elements along the reduction dimension with one e4m3 scale factor; by comparison, MXFP4 uses 𝐺=32 elements and one e8m0 scale factor. For a block 𝑥𝑔 ∈ R𝐺 , quantization and hardware reconstruction can be written as
𝑠𝑔 = 2
amax(𝑥𝑔 ) , 𝜌
𝑥ˆ𝑔 =
j 𝑥𝑔 m e2m1, 𝑠𝑔
𝑥˜𝑔 = 𝑠𝑔 𝑥ˆ𝑔 ,
(1)
SM
global memory Q, K, V
SS
shared memory Q, K, V tiles
by softmax, making quantized attention sensitive to outliers, normalization, and accumulation error. The SageAttention series [29–31] and HiFA4 [6] demonstrate that post-training quantization can accelerate attention without retraining. It can also consume low-bit KV caches directly [16], reducing memory use and avoiding dequantization.
Tensor Core tcgen05 MMA
2
TS
1
CUDA core, SFU softmax
tensor memory S, O, P
TMA
3 4
1 load
2 QK ⊤ and PV
3 softmax
4 store O
(a) FlashAttention4 data path
3
4 warps × 32
two query stages in flight fill all 512 columns
S0
S1
O0
Motivation
Low-bit attention can exploit the B200’s FP4 Tensor Cores while compressing the KV cache, reducing the compute and memory costs of long-context inference. However, Table 1 shows that existing systems target consumer Blackwell GPUs or NPUs, without addressing the B200’s tensor memory, operand-delivery paths, or serving-time decode.
O1
512 columns / CTA, 4 bytes per column
(b) Tensor memory architecture
3.1
Figure 2. Data path and tensor memory organization.
Blackwell’s gains concentrate in Tensor Cores: the B200 delivers 2.25 PFLOPS in BF16, compared with 1 PFLOPS on the H100, and 9 PFLOPS in FP4 [17, 19]. Exponential throughput and shared-memory bandwidth have not scaled proportionally [13, 18, 27], so reducing matmul precision yields no comparable reduction in softmax cost. The B200’s operanddelivery paths further amplify this asymmetry. Block-scaled FP4 MMA requires its operands to reside in shared memory and its scale factors to reside in tensor memory, whereas FP8 MMA can source its first operand directly from tensor memory. Moreover, BF16 FA4 double buffers the queries, while its two FP32 score and output accumulators, each with a tile size of 128, already occupy all 512 tensor memory columns available to each CTA, leaving limited capacity for additional scale factors. Precision selection is therefore jointly constrained by compute throughput, data movement, and contention for tensor memory capacity. We implement and measure four configurations on the B200: BF16 FA4, and three configurations that fix 𝑄𝐾 ⊤ to NVFP4 while using NVFP4, BF16, or FP8 for 𝑃𝑉 . Based only on Tensor Core MMA cycles, NVFP4 𝑃𝑉 should be fastest, followed by FP8 and then BF16, with theoretical speedups of 4.00×, 2.67×, and 1.60×, respectively. However, Figure 3(a) shows the opposite measured ordering. At 𝐿=64K, NVFP4 + FP8 is 1.38× faster than BF16 FA4, NVFP4 + BF16 achieves 1.11×, while NVFP4 + NVFP4 reaches only 0.87× and is slower than BF16 FA4. The result demonstrates that precision selection on the B200 cannot be based on Tensor Core peak throughput alone. Figure 3(b) decomposes the SM cycles per KV-block iteration and explains why the theoretical matrix-multiplication speedup does not translate into end-to-end operator speedup. In the BF16 configuration, the two matrix multiplications consume 1024 cycles, or 60% of total operator time. The BF16 FA4 implementation already approximates the exponential with a third-order polynomial on the CUDA cores to reduce the cost of the softmax exponential, and we adopt the same
where 𝜌=6 is the maximum magnitude of e2m1. In Equation 1, the Tensor Core performs the reconstruction 𝑥˜𝑔 = 𝑠𝑔 𝑥ˆ𝑔 internally during MMA. Block scaling limits quantization error locally, but introduces scale factors that must accompany the operands. We use NVFP4 because its 16-element blocks and scales better match attention activations than MXFP4 [30]. Different MMA precisions impose different operand-source and scale-factor requirements, making hardware details critical to both precision selection and performance bottlenecks. 2.2
FlashAttention and Quantization
Self-attention comprises two matrix multiplications and one softmax: √ 𝑆 = 𝑄𝐾 ⊤ / 𝑑, 𝑃 = softmax(𝑆), 𝑂 = 𝑃𝑉 (2) FlashAttention [3, 4] tiles 𝑄, 𝐾, and 𝑉 along the sequence dimension and processes KV blocks with online softmax, keeping intermediates on chip. √ Let the score for the 𝑡-th KV block be 𝑆 (𝑡 ) = 𝑄𝐾 (𝑡 )⊤ / 𝑑. FlashAttention maintains a row-wise maximum 𝑚, normalization factor ℓ, and output accumulator 𝑂: 𝑚 (𝑡 ) = max 𝑚 (𝑡 −1) , rowmax(𝑆 (𝑡 ) ) (3) 𝑃 (𝑡 ) = 𝑒 𝑆 ℓ (𝑡 ) = 𝑒
(𝑡 ) −𝑚 (𝑡 )
𝑚 (𝑡 −1) −𝑚 (𝑡 )
(4) (𝑡 )
ℓ (𝑡 −1) + rowsum 𝑃 (𝑡 −1) −𝑚 (𝑡 ) 𝑂 (𝑡 ) = diag 𝑒 𝑚 𝑂 (𝑡 −1) + 𝑃 (𝑡 ) 𝑉 (𝑡 )
Asymmetric Hardware Scaling
(5) (6)
Row maximum, exponentiation, row sum, and output rescaling in Equations 3–6 form the non-matmul softmax path. FA3 [23] hides its latency through warp specialization, overlapping a tile’s matmul with the preceding tile’s softmax. FA4 [27] reorganizes this pipeline for Blackwell tensor memory and fifth-generation Tensor Cores (Figure 2(a)). Quantized attention. Unlike static weights in linear-layer quantization, 𝑄, 𝐾, 𝑃, and 𝑉 are runtime activations coupled 3
Table 1. Positioning relative to the closest related systems. Work
Hardware
𝑄𝐾 ⊤ / 𝑃𝑉
Compensation
Partitioned decode
Workload
FA4 [27]
B200
BF16 / BF16
—
Incompatible with graph capture
Vision/language
FlashInfer [26]
B200
FP8 / FP8
—
Graph-compatible
Vision/language
SageAttn [31]
RTX 4090
INT8 / INT8
Softmax
—
Vision/language
SageAttn2 [29]
Hopper
INT4 / FP8
Softmax
—
Vision/language
SageAttn3 [30]
RTX 5090
NVFP4 / NVFP4
Softmax
—
Vision
HiFA4 [6]
Ascend NPU
HIF4 / HIF4
Softmax
—
Language
MpFA (ours)
B200
NVFP4 / FP8
MMA
Graph-compatible
Language
PV: NVFP4
PV: BF16
(a) 𝑄𝐾 ⊤ fixed to NVFP4
PV: FP8
MMA
non-MMA
500
1938
2000 1707
1545
1500
1235
1000 500
60% 41%
31%
13%
0
0
measured
BF16 +BF16
NVFP4 +NVFP4
NVFP4 +BF16
(b) Performance gap
NVFP4 +FP8
cycles per KV iteration
cycles per KV iteration
1
0.88
2
1.11
1.60
3
1.38
4
2.67
5
4.00
speedup vs BF16 FA4
roofline (MMA only)
417 MMA 384
400 300
256
200 100
93
85
row max
sum/rescale
0 exp2
others
(c) Performance breakdown
Figure 3. Four precision assignments for causal prefill at 𝐿=64K with GQA 32:8. technique. NVFP4 + FP8 reduces this component to 384 cycles, a 63% reduction, but total time falls only from 1707 to 1235 cycles, or 28%. Softmax, reductions, and data movement do not shrink with matrix-multiplication precision. Once attention enters the 4-bit regime, the dominant constraint therefore shifts from matrix multiplication to the softmax path. The reason NVFP4 + NVFP4 becomes slower is operand delivery rather than FP4 MMA throughput. If the probability matrix 𝑃, just generated in tensor memory by softmax, is used as an operand for block-scaled NVFP4 MMA, three additional costs arise. First, 𝑃 must be quantized by computing an absolute maximum over every 16 elements and generating scale factors; this work is placed on the already busy softmax path. Second, block-scaled MMA requires both operands from shared memory, so 𝑃 must leave tensor memory, be written to shared memory, and then be read by the matrix multiplication unit. Third, 𝑃’s scale factors need tensor memory capacity, competing with the score and output accumulators that already consume all 512 columns and requiring additional scale-factor transfers. The FP8 𝑃𝑉 path avoids all three costs. Without block scaling, FP8 MMA can read its first operand directly from tensor memory, allowing the probability matrix to remain in place. FP8 Tensor Cores also provide twice the throughput of BF16. Thus, an efficient B200 precision assignment
must consider both data format and operand-delivery path: NVFP4(𝑄𝐾 ⊤ )+FP8(𝑃𝑉 ) is the hardware-efficient point. 3.2
Inefficient Accuracy Compensation
Figure 3(c) further breaks down the execution time of NVFP4 + FP8. We obtain the softmax-stage costs through a compiletime ablation chain: we first rebuild the kernel with exp2 disabled, then additionally disable the row-max and rowsum/rescale stages, and compute successive differences. This procedure makes the stage costs additive. The softmax path consumes 595 cycles, or 48% of total operator time, with exp2 alone accounting for 417 cycles, exceeding the 384 cycles of the two quantized matrix multiplications. An 𝑀=𝑁 =128 KV block contains 16384 score elements. At 16 elements per SM cycle, fully serialized exponential evaluation would require 1024 cycles. The measured 417 cycles are lower because CUDA core polynomial evaluation overlaps with exponential unit execution, while matrix multiplication overlaps with the softmax path. Thus, once attention enters the 4-bit regime, non-matmul computation, particularly the softmax path, becomes the dominant bottleneck. Low-bit quantization still requires accuracy compensation. Prior methods typically smooth outliers by subtracting the means of 𝑄 and 𝐾 and adding a compensation factor back to the attention scores. The compensation must be 4
added before masking and before updating the running maximum; otherwise it breaks online-softmax normalization. Existing approaches consequently place additional loads and element-wise additions on the softmax bottleneck. In our measurements, this placement costs 18.6–22.7% of operator time, consuming more than half of the advantage of the unsmoothed low-bit attention. 3.3
or calibration. The remaining concern is the lower end of the FP8 range: in a forward traversal, an early attention sink can raise the running maximum and push ordinary-token probabilities into the e4m3 subnormal range. MpFA therefore traverses 𝐾 and 𝑉 in reverse KV-block order to delay the sink block. Since online softmax is traversal-order invariant, reversal changes only the finite-precision rounding path, following FA3 [23].
Insights from Profiling and Analysis 4.2
The analysis yields three principles for low-bit attention on data-center Blackwell GPUs. Principle 1: mixed precision. NVFP4(𝑄𝐾 ⊤ )+FP8(𝑃𝑉 ) accelerates 𝑄𝐾 ⊤ computation without quantizing 𝑃 to NVFP4, avoiding the associated data-movement and scale management overheads along the softmax path. Principle 2: reduce non-matmul overhead. Rank-one Tensor Core MMA compensation moves the cost from softmax to Tensor Cores, while a fine-grained pipeline improves utilization. Principle 3: specialize prefill and decode. Prefill processes the complete query sequence and is primarily constrained by the softmax path. Decode has only a small number of valid query rows at each step and is primarily constrained by insufficient parallelism and padded rows introduced by MMA atom.
4
NVFP4 block scaling assigns the dynamic range of each 16element block according to its amax. Stable channel outliers within a block can therefore directly reduce the effective precision of 4-bit quantization. Figure 4 shows activations collected from a running service: one head from one layer of each model over the first 256 tokens. The magnitudes of 𝑄 and 𝐾 form ridges along the channel dimension that persist across many tokens, while remaining approximately flat along the token dimension. Because NVFP4 blocks are partitioned along the channel axis, a small number of stable outlier channels can consume most of the representable range of their blocks. These outliers are determined by fixed model parameters rather than transient input tokens, making them suitable for smoothing. In contrast, 𝑉 varies comparably across the channel and token dimensions and is therefore not smoothed, consistent with prior work [29]. As 𝑉 is not centered, it requires no mean recovery and is unaffected by rank-one compensation. It uses FP8 e4m3 (three mantissa bits versus one for NVFP4) with FP32 accumulation. Since Í 𝑂 = 𝑗 𝑃 𝑗 𝑉𝑗 is a convex combination, rounding errors in 𝑉 enter the output as a weighted average rather than a sum that grows linearly with sequence length. Sequence centering. To mitigate the stable channel outlier structure, we subtract a sequence-level mean for each layer and head before quantization. For the set of valid tokens C, we define 1 ∑︁ 1 ∑︁ 𝐾 [𝑡, ℎ, :], 𝑄¯ [ℎ, :] = 𝑄 [𝑡, ℎ, :] (7) 𝐾¯ [ℎ, :] = |C| 𝑡 ∈ C |C| 𝑡 ∈ C
Quantization Algorithm
Following the first two principles in §3, MpFA uses NVFP4 for 𝑄𝐾 ⊤ , FP8 for 𝑃𝑉 , and sequence centering with BF16 rank-one Tensor Core MMA to compensate for channeloutlier-induced NVFP4 error in 𝑄 and 𝐾. 4.1
QK4PV8 Quantization
MpFA quantizes 𝑄 and 𝐾 to NVFP4 with one e4m3 scale factor per 16 head-dimension elements, and quantizes 𝑃 and 𝑉 to FP8. The online softmax states 𝑚, ℓ, and the 𝑃𝑉 accumulator remain FP32, and the operator returns BF16. This assignment reduces 𝑄𝐾 ⊤ computation and KV storage without the block-scale management required by fully NVFP4 𝑃𝑉 . Quantizing 𝑃 does not require a dynamically computed scale factor for each block, but should account for the limited range of FP8 e4m3 and the numerical degradation caused by attention sinks. From Equation 4, each unnormalized probability generated by online softmax satisfies (𝑡 )
Outlier Structure and Sequence Smoothing
The system obtains these two statistics during prefill and quantizes the centered values ¯ 𝑄𝑐 = 𝑄 − 𝑄, 𝐾𝑐 = 𝐾 − 𝐾¯ . We partition 𝑄𝑐 , 𝐾𝑐 , and 𝑉 along the token dimension into {𝑄𝑐,𝑖 }, {𝐾𝑐,𝑗 }, and {𝑉𝑗 }, with quantized blocks 𝑄ˆ 𝑐,𝑖 , 𝐾ˆ𝑐,𝑗 , and 𝑉ˆ𝑗 and scales 𝑠𝑄,𝑖 and 𝑠𝐾,𝑗 . The removed means become recoverable compensation factors, while centering is fused into the quantization write path.
(𝑡 )
0 < 𝑃𝑖(𝑡𝑗 ) = 𝑒 𝑆𝑖 𝑗 −𝑚𝑖 ≤ 1. MpFA therefore uses a fixed scale factor 𝑠𝑃 = 448 for 𝑃: 𝑃b𝑖(𝑡𝑗 ) = caste4m3 𝑠𝑃 𝑃𝑖(𝑡𝑗 ) /𝑠𝑃 .
4.3
Rank-One Tensor Core Compensation
When both 𝑄 and 𝐾 are centered, the removed mean terms must be restored before softmax. Expanding 𝑆 = 𝑄𝐾 ⊤ and defining a key-dependent compensation vector gives ¯ 𝑐⊤ + 𝑄 𝐾¯ ⊤ 1⊤ 𝑆 = 𝑄𝑐 𝐾𝑐⊤ + 1𝑀 𝑄𝐾 (8)
Because 448 is the largest normal e4m3 value and 𝑃𝑖(𝑡𝑗 ) ≤ 1, this analytic scale applies across layers, heads, and inputs without runtime amax statistics, per-block scale generation,
𝑁
5
4
0
cha
50
100
nne
l
6 4
7.5 5.0
2
5
2
2.5
0
0
0
0.0
200 100 n 0
15 10
ke
to
(a) Llama-3.1-8B |𝑄 |
0
cha
50
100
nne l
200 100 n 0
0
ke to
50
100 cha nne l
(b) Llama-3.1-8B |𝐾 |
200 100 n 0
ke
to
(c) Qwen3-14B |𝑄 |
0
cha
50
100
nne
200 100 n 0
l
ke
to
(d) Qwen3-14B |𝐾 |
Figure 4. Distributions of |𝑄 | and |𝐾 | for one head at layer 16 of Llama-3.1-8B and layer 20 of Qwen3-14B. Δ𝑆 [ 𝑗] ≜ 𝑄¯ · 𝐾𝑐 [ 𝑗, :] ⊤
Algorithm 1 Algorithm design of MpFA. Δ𝑆,𝑗 is the com√ pensation vector. The 1/ 𝑑 factor is folded into the 𝜆.
(9)
Here Δ𝑆,𝑗 denotes the entries of Δ𝑆 [ 𝑗] in KV block 𝑗; GQA uses one Δ𝑆 vector per query head. The third term is constant across keys within each softmax row and cancels because softmax(𝑥 + 𝑐1) = softmax(𝑥). Only the key-dependent second term must be restored after centering: 𝑆ˆ = deq 𝑄ˆ 𝑐 𝐾ˆ𝑐⊤ + 1𝑀 Δ𝑆 (10)
Input: 𝑄, 𝐾, 𝑉 in BF16; block sizes 𝐵𝑞 , 𝐵𝑘𝑣 ; 𝑇𝑚 = ⌈𝑁𝑞 /𝐵𝑞 ⌉ query blocks and 𝑇𝑛 = ⌈𝑁𝑘𝑣 /𝐵𝑘𝑣 ⌉ KV blocks; 𝑀 = 𝐵𝑞 Output: 𝑂 in BF16 1: 𝑄¯ ← mean(𝑄), 𝐾¯ ← mean(𝐾) ⊲ outside kernel ¯ 𝐾𝑐 ← 𝐾 − 𝐾¯ 2: 𝑄𝑐 ← 𝑄 − 𝑄, 3: (𝑠𝑄 , 𝑄ˆ 𝑐 ) ← QuantNVFP4 (𝑄𝑐 ); (𝑠𝐾 , 𝐾ˆ𝑐 ) ← QuantNVFP4 (𝐾𝑐 ) 4: 𝑉ˆ ← QuantFP8 (𝑉 ) ⊲ fused into KV-cache write ¯ 𝐾𝑐⊤ 5: Δ𝑆 ← GEMVbf16 𝑄, ⊲ Equation 9 6: for 𝑖 = 1 to 𝑇𝑚 do ⊲ kernel main loop 7: 𝑚𝑖,0 ← −∞, ℓ𝑖,0 ← 0, 𝑂𝑖,0 ← 0 8: for 𝑗 = 1 to 𝑇𝑛 do 9: load 𝑄ˆ 𝑐,𝑖 , 𝑠𝑄,𝑖 , 𝐾ˆ𝑐,𝑗 , 𝑠𝐾,𝑗 , 𝑉ˆ𝑗 , Δ𝑆,𝑗 ⊲ Equation 11 ˆ 𝑐,𝑖 , 𝑠𝑄,𝑖 , 𝐾ˆ𝑐,𝑗 , 𝑠𝐾,𝑗 ) 10: 𝑆𝑖,𝑗 ← MMAbf16 (1𝑀 , Δ𝑆,𝑗 ) + MMA ( 𝑄 fp4 11: 𝑚𝑖,𝑗 ← max 𝑚𝑖,𝑗 −1, rowmax(𝑆𝑖,𝑗 ) √ √ 12: 𝑃𝑖,𝑗 ← 2𝜆 (𝑆𝑖,𝑗 −𝑚𝑖,𝑗 ) , 𝜆 = log2𝑒/ 𝑑 ⊲ 1/ 𝑑 folded here 13: ℓ𝑖,𝑗 ← 2𝜆 (𝑚𝑖,𝑗 −1 −𝑚𝑖,𝑗 ) ℓ𝑖,𝑗 −1 + rowsum(𝑃𝑖,𝑗 ) ⊲ FP32 𝑃𝑖,𝑗 14: 𝑃ˆ𝑖,𝑗 ← Quant FP8 (𝑃𝑖,𝑗 ) 15: 𝑂𝑖,𝑗 ← diag 2𝜆 (𝑚𝑖,𝑗 −1 −𝑚𝑖,𝑗 ) 𝑂𝑖,𝑗 −1 + MMAfp8 (𝑃ˆ𝑖,𝑗 , 𝑉ˆ𝑗 )
Here Δ𝑆 is added before masking and before updating the running maximum. This ordering prevents compensation from reaching masked positions and keeps the running maximum consistent with the normalized scores. Rank-one MMA implementation. The compensation term in Equation 10 is the outer product 1𝑀 Δ𝑆 within a score tile. Rather than adding it element-wise after computing 𝑄ˆ 𝑐 𝐾ˆ𝑐⊤ , we first issue a BF16 Tensor Core MMA instruction on the same score accumulator: 𝑆 acc ← 1𝑀 Δ𝑆 , 𝑆 acc += deq 𝑄ˆ 𝑐 𝐾ˆ𝑐⊤ . (11)
16: end for 17: 𝑂𝑖 ← diag(ℓ𝑖,𝑇𝑛 ) −1𝑂𝑖,𝑇𝑛 18: end for 𝑇𝑚 19: return 𝑂 = {𝑂𝑖 }𝑖=1
For the BF16 MMA reduction granularity 𝐾=16, the constant 𝐴 operand has 1 only in its first column, with the remaining columns and corresponding Δ𝑆 components zeroed. The accumulation is therefore exactly 1𝑀 Δ𝑆 , moving compensation from the softmax path to otherwise idle Tensor Cores without changing online softmax. Compared with element-wise addition, this introduces one BF16 rounding with relative error at most 2−8 .
5
This work is fused with the inference framework’s KV-cache write path (§6), keeping its overhead low. The second stage is the attention main loop, which consumes only the quantized data and compensation factors. In the algorithm, 𝑖 and 𝑗 index query blocks and KV blocks, respectively. 𝑄ˆ 𝑐,𝑖 is the 𝑖-th quantized centered query block; 𝐾ˆ𝑐,𝑗 and 𝑉ˆ𝑗 form the 𝑗-th KV block; and 𝑠𝑄,𝑖 and 𝑠𝐾,𝑗 are their NVFP4 block scale factors. 𝑆𝑖,𝑗 and 𝑂𝑖,𝑗 denote the score and output states after processing KV block 𝑗, while 𝑚𝑖,𝑗 and ℓ𝑖,𝑗 are the running maximum and normalization factor of online softmax. The only quantization inside the loop converts 𝑃𝑖,𝑗 to FP8. Since its range has a known upper bound (§4.1), this conversion uses a fixed scale factor.
High Performance QK4PV8 FA Kernel
This section presents MpFA’s unified QK4PV8 computation flow, followed by tensor memory reuse, rank-one compensation placement, and decode-specific optimizations. 5.1
Fine-Grained Pipeline
Algorithm 1 describes MpFA’s two-stage execution flow. The first stage runs outside the attention operator: it computes ¯ centers 𝑄 and 𝐾, quantizes 𝑄𝑐 and 𝐾𝑐 to NVFP4, 𝑄¯ and 𝐾, quantizes 𝑉 to FP8, and computes the per-token compensation vector Δ𝑆 with one GEMV according to Equation 9. 6
KV block j TMA load
K̂ c NVFP4 + sK · V ̂ FP8
MMA
seed Q̂c Kĉ ⊤ 0 ΔS
sQ , sK copy tcgen05.cp
s (0) →S (1)
tensor memory 512 columns
s (1)
P̂ V ̂ 0
j+1 ΔS, j
Q buf 0
seed Q̂c Kĉ ⊤ 1 ΔS
P̂ V ̂ 1
s (1) →S (0) P̂ (0) 64 col S (0) 128 col
K̂ c NVFP4 + sK · V ̂ FP8
seed Q̂c Kĉ ⊤ 0 ΔS s (0) →S (1) P̂ (1) 64 col S (1) 128 col
s (0)
P̂ V ̂ 0
j+2 Q buf 1
seed Q̂c Kĉ ⊤ 1 ΔS
ΔS, j P̂ V ̂ 1
s (1) →S (0)
K̂ c NVFP4 + sK · V ̂ FP8
seed Q̂c Kĉ ⊤ 0 ΔS
P̂ V ̂ 0
s (0) →S (1)
O (0) 128 col
seed Q̂c Kĉ ⊤ 1 ΔS
O (1) 128 col
stage 0: mask · exp · Σ · → FP8
stage 0: mask · exp · Σ · → FP8
⋯
softmax WG1
stage 1: mask · exp · Σ · → FP8
stage 1: mask · exp · Σ · → FP8
⋯
correction
rescale O by e
rescale O by e
P̂ V ̂ 1
s (1) →S (0)
softmax WG0
mj − 1 − mj
ΔS, j
Q buf 0
mj − 1 − mj
epilogue
O/ℓ O → BF16
softmax drains S, writes P̂ into its upper half s (r) borrows stage 1−r's S columns
QK 0/1 write S (0) /S (1) , FP8 PV reads P̂ back correction rescales O in place
Figure 5. MpFA’s steady-state pipeline. The seven execution channels for five warp roles within one CTA advance over consecutive KV blocks, together with their tensor memory accesses. Warp specialization. MpFA retains the five warp roles of FA4 [27], but adapts them for QK4PV8: load, MMA, softmax, correction, and epilogue, as shown in Figure 5. The load warp issues asynchronous copies of 𝐾ˆ𝑐,𝑗 , 𝑠𝐾,𝑗 , 𝑉ˆ𝑗 , the current query block and its scale factor, and Δ𝑆,𝑗 from global memory to on-chip memory. The MMA warp issues two MMAs for the current score state: a rank-one BF16 MMA first writes Δ𝑆,𝑗 , followed by a block-scaled NVFP4 MMA that accumulates ⊤ ; an FP8 MMA then computes 𝑃ˆ 𝑉ˆ . The two softmax 𝑄ˆ 𝑐,𝑖 𝐾ˆ𝑐,𝑗 𝑖,𝑗 𝑗 warpgroups apply the mask, update 𝑚𝑖,𝑗 and ℓ𝑖,𝑗 , and convert 𝑃𝑖,𝑗 to 𝑃ˆ𝑖,𝑗 in place. The correction warpgroup rescales the output state and performs final normalization, while the epilogue warp writes the BF16 output back to global memory. We double-buffer the query blocks (𝑄𝑐 ) to overlap MMA with softmax. One CTA processes two query stages, indexed by 𝑟 ∈ {0, 1}. Each stage owns one 𝑆 (𝑟 ) and one 𝑂 (𝑟 ) , corresponding to an (𝑆𝑖,𝑗 , 𝑂𝑖,𝑗 ) state in the algorithm. The probability matrix 𝑃ˆ𝑖,𝑗 receives no separate tensor-memory allocation; it is written into the upper half of the current 𝑆 (𝑟 ) region, and the FP8 𝑃𝑉 MMA reads it back from that region. Thus, stage 𝑟 ’s NVFP4 𝑄𝐾 ⊤ writes only 𝑆 (𝑟 ) , its softmax reads 𝑆 (𝑟 ) and writes 𝑃ˆ (𝑟 ) in the upper half, and its FP8 𝑃𝑉 MMA reads 𝑃ˆ (𝑟 ) from the same region. The two stages alternate between 𝑆 (0) /𝑃ˆ (0) and 𝑆 (1) /𝑃ˆ (1) , allowing one stage’s softmax to overlap with the other stage’s matrix multiplication. This double buffer hides softmax latency and enables reuse of the scale factors described next. Asynchronous pipeline for scale factors in tensor memory. Each CTA is allocated 512 tensor memory columns. The score and output states of the two query stages require 128 columns each, so the four accumulators consume the entire capacity. NVFP4 scale factors for 𝑄 and 𝐾 cannot receive
independent allocations. MpFA resolves this conflict by exploiting asymmetric lifetimes: 𝑆 (𝑟 ) and 𝑂 (𝑟 ) must remain resident through a long softmax or main-loop interval, whereas 𝑠𝑄,𝑖 and 𝑠𝐾,𝑗 are read only when the corresponding NVFP4 MMA is issued. We asynchronously place FP8 scale factors in reused storage. The scale factors needed by stage 𝑟 are not placed on its own 𝑆 (𝑟 ) region; instead, they borrow columns from 𝑆 (1−𝑟 ) : 𝑠 (0) occupies the starting columns of 𝑆 (1) , and 𝑠 (1) occupies those of 𝑆 (0) (the dashed cells at the left of the score regions in Figure 5). The double-buffer timing makes this safe. While stage 𝑟 issues 𝑄𝐾 ⊤ , stage 1 − 𝑟 ’s score region is in the gap after softmax and 𝑃𝑉 have consumed it but before the next MMA overwrites it. In contrast, 𝑆 (𝑟 ) is the target of the current MMA and cannot simultaneously hold its scale factors. Scale-factor reads and writes form an independent pipeline with a phase distinct from data loading and matrix multiplication. Each query stage owns a set of tcgen05.cp copies and its own barrier phase (the two 𝑠𝑄 , 𝑠𝐾 copy boxes in Figure 5). Copies are issued early enough for the factors to arrive before the corresponding MMA; their transfer overlaps with the other stage’s softmax and matrix multiplication and therefore does not enter the critical path. Correctness follows from two phase conditions: scale factors must arrive before their MMA is issued, and a new write to a score region can occur only after that region has been consumed by softmax and the 𝑃𝑉 MMA. The probability matrix 𝑃 uses the same lifetime reuse but occupies only the upper half of the score region, requiring no additional tensor-memory allocation. This allows scale factors, attention scores, and outputs to remain resident within the fixed 512-column budget. Other attention optimizations. MpFA builds the kernel on the CUTLASS and CuTe programming primitives. For causal attention, MpFA balances the triangular workload 7
using longest-processing-time-first scheduling and overlaps TMA and MMA with a ping-pong KV pipeline. It also distributes softmax exponentiation between a polynomial approximation and hardware 2𝑥 instructions to balance the FMA and SFU pipelines. Moreover, attention scaling is fused into exponentiation: 𝑄 and 𝐾 remain in the unscaled inner√ product domain, and the softmax constant 𝜆 = log2𝑒/ 𝑑 is folded into the exponent (2𝜆 (𝑆 −𝑚) = 2𝜆𝑆 −𝜆𝑚 ), avoiding separate scaling of the scores and compensation term Δ𝑆,𝑗 . 5.2
following prior work [20]. The backend supports paged prefill and decode, GQA, masking, and CUDA Graph capture, while SGLang handles the remaining model execution. During decode, fixed-size cache pages and preallocated tensors keep the number of KV partitions constant after graph capture. MpFA returns BF16 outputs through the standard attention interface. With smoothing, the per-token compensation Δ𝑆 follows the KV cache’s physical page-table indices and resides in a side cache directly accessed during decode. Online quantization. MpFA fuses NVFP4 quantization of 𝑄, low-bit 𝐾/𝑉 cache writes, and compensation generation into SGLang’s cache-write path. It stores 𝐾 in NVFP4 with one e4m3 scale byte per 16 elements and 𝑉 in FP8, allowing direct consumption without dequantization. For equaldimensional 𝐾 and 𝑉 , the layout uses 1/2.56 of the BF16 KV-cache footprint: 1/2 + 1/16 bytes per 𝐾 element and one byte per 𝑉 element, versus two bytes per element for each BF16 tensor. With smoothing, prefill computes perlayer, per-head 𝑄¯ and 𝐾¯ over valid tokens and freezes them for the sequence lifetime. Decode centers only the newly appended token against these per-request means without recomputing them. This avoids degenerate single-token query centering that would shift the entire score correction to Δ𝑆 . For the new token, centering and compensation generation are fused into the cache-write kernel. Without smoothing, no side cache is allocated. Kernel measurements exclude quantization statistics and compensation generation; end-to-end measurements include quantization, sequence centering, and KV-cache writes.
MMA Compensation
For each (𝑖, 𝑗), the MMA warp first issues MMAbf16 (1𝑀 , Δ𝑆,𝑗 ), then accumulates the NVFP4 MMAnvfp4 (𝑄ˆ 𝑐,𝑖 , 𝑠𝑄,𝑖 , 𝐾ˆ𝑐,𝑗 , 𝑠𝐾,𝑗 ) into the same score state 𝑆𝑖,𝑗 . Because the compensation has already been written to the score accumulator before softmax, the softmax warp reads complete scores without an additional element-wise addition or a change to the onlinesoftmax instruction sequence. The compensation vector Δ𝑆,𝑗 is much smaller than 𝐾ˆ𝑐,𝑗 and 𝑉ˆ𝑗 . The loading warp prefetches it and publishes it with the same barrier phase as the current KV block. Rank-one MMA and NVFP4 𝑄𝐾 ⊤ consequently share one dependency and require no additional synchronization. Relative to the same build without smoothing, the prefill attention operator adds only 1.3–2.5%; placing compensation on the softmax warp costs 18.6–22.7% (§7.4). 5.3
Decode-Specific Optimizations
Decode uses the same representations and MMA precisions as prefill but is parallelism-bound. MpFA packs GQA heads, partitions the KV dimension, removes softmax work from padded rows, and merges the resulting local states. GQA lets 𝐻𝑞 /𝐻𝑘𝑣 query heads share one KV head, so MpFA packs each group into the MMA 𝑀 dimension and assigns one (request, KV-head) pair to each CTA. It partitions the history into fixed 8 × 128-token chunks; each CTA maintains local (𝑚, ℓ, 𝑂) states in preallocated workspaces, making the path CUDA Graph capturable and increasing parallelism over split-KV-disabled BF16 FA4. For GQA 32:8, FP4 MMA fixes 𝑀 = 128 although each group has only four real rows. MpFA replicates these rows and processes their corresponding columns as four diagonal 32 × 32 tiles, reducing softmax work from 32 × 128 to 32 × 32 per row range and the full-tile 2𝑥 evaluations from 128 × 128 to 128 × 32. The four warps produce four states per partition, yielding 4𝑁𝑝 states merged by log-sum-exp. Eight independent state streams and shared-memory scale prefetching hide the reduction’s memory latency; the same strategy extends to other model shapes.
6
7
Evaluation
7.1
Experimental Setup
Hardware and software. We conduct all experiments on a single NVIDIA B200 using CUDA 13.0, PyTorch 2.11.0, and SGLang 𝑣0.5.15. For end-to-end experiments, we enable CUDA Graph capture and disable Radix Cache [32]. Models and accuracy tasks. We evaluate Llama-3.1-8B and Qwen3-14B, covering GQA 32:8 and 40:8; both have 128dimensional heads. We extend Qwen3-14B from its native 40K to 128K using YaRN, while Llama-3.1-8B uses its native 128K context. This extension may affect accuracy but not performance measurements. We evaluate MMLU [11] (5-shot, all 57 subjects), HellaSwag [28] (20-shot), WinoGrande [22] (0-shot), GSM8K [2] (5-shot with generated reasoning), and RULER-QA [12], which pads SQuAD and HotpotQA passages to 64K tokens. Each backend uses the same examples and prompts with greedy decoding (𝑡𝑒𝑚𝑝𝑒𝑟𝑎𝑡𝑢𝑟𝑒 = 0). Baselines. We compare BF16 FA4, FP8 FlashInfer, and three precision assignments for 𝑄𝐾 ⊤ and 𝑃𝑉 : NVFP4+NVFP4, NVFP4+BF16, and NVFP4+FP8. MpFA is the complete QK4PV8 design with rank-one smoothing. For decode kernel tests, we tune BF16 FA4’s split count per context length and use
LLM Serving System
Serving system. To evaluate end-to-end performance, we integrate MpFA into SGLang [32] as an attention backend, 8
2.92× 3.13× 3.12×
2.98× 3.46× 3.45×
10
2.25× 2.81× 2.60×
20
2.15× 2.48× 2.45×
30
1.14× 1.26× 1.14×
MpFA(smoothing)
1.36× 1.43× 1.29×
NVFP4+BF16 NVFP4+FP8
TFLOPS
1.17× 0.91× 1.11× 1.39× 1.34×
FlashInfer FP8 NVFP4+NVFP4
1.16× 0.88× 1.11× 1.38× 1.33×
1.13× 0.82× 1.09× 1.37× 1.31×
1.06× 0.73× 1.05× 1.27× 1.26×
1000
0.64× 0.43× 0.66× 0.76× 0.73×
TFLOPS
2000
0.90× 0.61× 0.96× 1.12× 1.09×
BF16 FA4 BF16 FA4 + split-KV (tuned)
0
0 4K
8K
16K
32K
64K
4K
128K
8K
16K
32K
64K
128K
sequence length
sequence length
(a) causal prefill
(b) decode
Figure 6. Attention-kernel throughput on a B200 with GQA 32:8 and ℎ𝑒𝑎𝑑_𝑑𝑖𝑚=128, the shape of Llama-3.1-8B. Speedups above the bars use BF16 FA4 in (a) and split-KV-enabled BF16 FA4 in (b).
1.65× 3.87× 3.53×
20
4K
8K
1.64× 3.18× 3.17×
1.34× 1.43× 1.30×
30
1.64× 2.55× 2.39×
40
10
500
MpFA(smoothing) 1.11× 1.25× 1.14×
NVFP4+BF16 NVFP4+FP8
2.43× 2.94× 2.66×
1.18× 0.92× 1.12× 1.39× 1.35×
1.17× 0.89× 1.12× 1.38× 1.34×
1.15× 0.84× 1.11× 1.37× 1.33×
FlashInfer FP8 NVFP4+NVFP4
TFLOPS
1000
1.11× 0.76× 1.08× 1.34× 1.29×
1500
0.77× 0.48× 0.75× 0.87× 0.83×
TFLOPS
2000
0.98× 0.63× 1.00× 1.17× 1.14×
BF16 FA4 BF16 FA4 + split-KV (tuned)
0
0 4K
8K
16K
32K
64K
128K
16K
32K
64K
128K
sequence length
sequence length
(a) causal prefill
(b) decode
Figure 7. Attention-kernel throughput for Qwen3-14B (GQA 40:8). The baselines and speedup denominators match Figure 6. > BF16 > NVFP4 + NVFP4, consistent with §3.1. Across 16K–128K, MpFA sustains 1672–1817 TFLOPS and achieves a 1.31× geometric-mean speedup over BF16 FA4; the direct variant reaches 1.35×, showing that smoothing preserves most of the gain. FP8 FlashInfer reaches 1.12× over BF16, while MpFA is 1.17× faster than FlashInfer. The benefit follows from reducing the compute-intensive 𝑄𝐾 ⊤ from FP8 to NVFP4. Qwen3-14B’s GQA 40:8 shape shows the same trend (Figure 7), with 1.33×/1.15× speedups over BF16 FA4/FP8 FlashInfer. Fixed overhead limits short sequences: at 4K, MpFA and its direct variant reach only 0.74× and 0.76× of BF16, versus 1.09× and 1.12× at 8K. The crossover therefore falls between 4K and 8K, below which matrix multiplication does not dominate. MpFA is thus suited to end-to-end inference at sequence lengths beyond 8K. NVFP4 + BF16 reaches 1.05– 1.11× at ≥ 16K, confirming the benefit of quantizing 𝑄𝐾 ⊤ alone. Relative to FP8, however, its 𝑃𝑉 MMA and sharedmemory traffic for 𝑉 both double, and it cannot use a low-bit 𝑉 cache. Fully NVFP4 is slowest from 4K onward, reaching only 0.87–0.90× BF16 at long sequences because of the additional 𝑃𝑉 execution-path costs identified in §3.1. Decode workloads. At 128K, MpFA and its direct variant reach only 22.8 and 25.1 TFLOPS, respectively, confirming that long-context decode is parallelism-bound rather than
FlashInfer’s default split-KV policy. MpFA instead uses fixed KV-partition and column-split configurations. For end-toend tests, BF16 FA4 runs one CTA per request because its split-KV path preallocates workspaces per num_splits and cannot be captured by CUDA Graph. FlashInfer and MpFA use graph-compatible split-KV, with FlashInfer using an FP8 KV cache. Integrated as an SGLang attention backend, MpFA supports the same production features as the baselines: continuous batching, paged KV cache, CUDA Graph capture, and head_dim (128 for both models). We measure prefill/decode kernel throughput, time to first token (TTFT), time per output token (TPOT), output-token throughput (OTPS) across concurrency levels, end-to-end accuracy, and ablations. Following prior work [5], we report the mean of five performance runs. 7.2
FlashAttention Kernel Performance
Figure 6 compares causal prefill and decode throughput across precision configurations, using BF16 FA4 and split-KVenabled BF16 FA4 as their respective denominators. NVFP4 + FP8 is the hardware-efficient prefill assignment; decode also depends on KV partitioning and softmax scheduling. MpFA includes smoothing, while direct NVFP4 + FP8 isolates its compensation cost. Prefill workloads. At 16K and above, measured throughput follows direct NVFP4 + FP8 > MpFA > NVFP4 + BF16 9
Llama-3.1-8B, FlashInfer FP8 Qwen3-14B, FlashInfer FP8
Llama-3.1-8B, MpFA(direct) Qwen3-14B, MpFA(direct)
1.00
speedup
speedup
1.25
0.75 0.50 0.25 4K
8K
16K
32K
64K
Llama-3.1-8B, MpFA(smoothing) Qwen3-14B, MpFA(smoothing)
7 6 5 4 3 2 1 0 4K
128K
8K
16K
32K
64K
128K
input length
input length
(a) Time to first token (TTFT)
(b) Time per output token (TPOT)
Figure 8. End-to-end latency speedup over BF16 FA4 at concurrency 1. MpFA(smoothing)
Tensor-Core-bound. We compare against three baselines with distinct deployment constraints. The first is production-deployable BF16 FA4. Its split-KV path preallocates workspaces according to num_splits and cannot be captured by CUDA Graph, so SGLang uses one CTA to process the full history sequentially. Relative to this configuration, MpFA improves throughput from 2.63× at 4K to 13.04× at 128K; the direct variant reaches 2.83× and 14.37×, respectively. The second is BF16 FA4 with split-KV enabled. We sweep num_splits∈ {1, . . . , 32} and select the best value at each length; the optimum is 16 or 32 at ≥ 8K. Against this perlength-tuned baseline, MpFA achieves a 2.08× average speedup across 8K–128K without per-length tuning, compared with 2.17× for the direct variant. MpFA peaks at 3.45× at 8K and remains 1.14× faster at 128K, due to GQA packing, fixed KV partitions, column splitting, and parallel state merging. The speedup decreases with sequence length: partitioning exposes substantially more parallelism than the baseline at short-to-medium lengths, whereas longer KV histories make both systems HBM-bandwidth-bound, limiting the benefit of NVFP4 compute. The third is FlashInfer’s FP8 Blackwell FlashAttention, which supports split-KV under CUDA Graph. MpFA achieves a 1.06× geometric-mean speedup across 8K–128K, leading by 1.11–1.14× through 32K but falling to 0.95× at 64K and near parity (0.99×) at 128K. The direct variant reaches 1.11× in geometric mean and remains 1.05×/1.09× faster at 64K/128K. The gap reflects smoothing: the Δ𝑆 side cache adds about 4% KV traffic, which is mostly absorbed by L2 through 32K but reaches HBM beyond 64K, making MpFA about 10% slower than the direct variant. The modest margin over FlashInfer also reflects the decode regime: at concurrency 1, reading the KV history dominates, so NVFP4’s compute advantage over FP8 is largely unused. Moreover, block-scaled FP4 MMA requires 𝑀=128, while a GQA group has only 𝐻𝑞 /𝐻𝑘𝑣 real query rows; row replication and column splitting cannot eliminate all padded work. These decode results generalize to GQA 40:8 (Figure 7): MpFA peaks at 3.53× over split-KV BF16 FA4 at 8K and stays 1.14× at 128K. Its larger query-to-KV
FlashInfer FP8
conc. 1
16K
32K
conc. 8
conc. 16
OTPS speedup
3.5 3.0 2.5 2.0 1.5 1.0 4K
8K
64K
128K
input length
Figure 9. Output throughput of SGLang serving relative to BF16 FA4 for Qwen3-14B. ratio feeds more compute per KV read, yielding 0.97×/1.02× of FP8 FlashInfer at 64K/128K. 7.3
End-to-End Inference Performance
Latency. Figure 8 compares MpFA, its direct variant, BF16 FA4, and FP8 FlashInfer under identical serving parameters. Across both models and 16K–128K, MpFA improves TTFT and TPOT over BF16 by 1.11× and 3.30× in geometric mean. At 128K, TPOT falls from 42.4 to 6.5 ms for Llama-3.18B (6.58×) and from 53.8 to 9.5 ms for Qwen3-14B (5.67×), slightly exceeding FlashInfer’s 6.38×/5.59×. Prefill gains are smaller because non-attention computation remains unchanged: at 128K, MpFA improves TTFT by 1.16×/1.14×. Although FP8 FlashInfer outperforms BF16 FA4 (Figure 6), separately producing its FP8 KV cache offsets this kernel gain end to end, giving 0.91–0.96× BF16 TTFT. MpFA fuses sequence centering, quantization, and cache writes (Figure 11), improving TTFT by up to 1.16× over BF16 FA4 and 1.16–1.27× over FP8 FlashInfer at 32K–128K. Fixed overhead limits both latency metrics at short sequences, so we make no stable speedup claim for inputs of ≤ 2K. Relative to MpFA (direct), smoothing adds 1.5–7.3% to TTFT and 3.0–6.2% to TPOT over 4K–128K; reduced KV traffic dominates as context grows. Throughput. Figure 9 reports output throughput relative to end-to-end BF16 FA4 with fixed input length and 1024 output tokens. At concurrency 1, MpFA achieves a 2.81× geometric-mean speedup across both models and 16K–128K, 10
Table 2. End-to-end accuracy. Rec. is the fraction of NVFP4+FP8 loss (vs. BF16) that MpFA recovers. Llama-3.1-8B
Qwen3-14B
Dataset
BF16
FP8
NVFP4+FP8
MpFA
Rec.
BF16
FP8
NVFP4+FP8
MpFA
Rec.
MMLU
0.685
0.683
0.644 (−0.041)
0.674 (−0.011)
73%
0.787
0.787
0.776 (−0.011)
0.785 (−0.002)
82%
HellaSwag
0.785
0.785
0.765 (−0.020)
0.780 (−0.005)
75%
0.774
0.774
0.768 (−0.006)
0.773 (−0.001)
83%
WinoGrande
0.752
0.751
0.735 (−0.017)
0.745 (−0.007)
59%
0.727
0.726
0.712 (−0.015)
0.720 (−0.007)
53%
GSM8K
0.761
0.757
0.653 (−0.108)
0.736 (−0.025)
77%
0.887
0.882
0.877 (−0.010)
0.882 (−0.005)
50%
RULER-QA
0.624
0.619
0.533 (−0.091)
0.578 (−0.047)
49%
0.656
0.656
0.592 (−0.064)
0.607 (−0.049)
24%
2.3
5
2.5
32K
64K
4.55 0.45
0.98
10−1
0.12
10
100
0.27
15
101
0.02
added time (ms, log)
18.9
20
18.9
25
1.3
added time (% of kernel)
22.7
on the softmax warps (prior work) on the Tensor Core (ours)
reaching 4.47×/3.76× at 128K (Llama-3.1-8B/Qwen3-14B). Absolute throughput then increases from 21.2 to 94.8 tok/s and from 16.2 to 61.0 tok/s. The smaller throughput gain than TPOT reflects the more modest 1.14–1.16× TTFT speedup. This speedup is measured against the production-deployable BF16 FA4, whose split-KV path cannot be graph-captured. However, the controlled kernel comparison with per-length split-KV tuning (Figure 6b) shows that MpFA remains faster at every length (2.08× on average and 1.14× at 128K), confirming that split-KV does not change the end-to-end performance ordering. Relative to FP8 FlashInfer, MpFA leads across 32K–128K by a 1.06× geometric mean, reaching 1.11–1.12× at 128K on both models. The margin grows with context as sequence centering, quantization, and cache writes amortize and MpFA’s lower KV-cache traffic comes to dominate. Higher concurrency narrows MpFA’s advantage over BF16 but preserves the long-context benefit: at concurrency 8, it achieves a 1.33× average speedup across both models and 8K–128K, reaching 1.70×/1.89× at 128K. For Qwen3-14B at concurrency 16, it overtakes BF16 at 16K (1.04×) and reaches 1.13× at the fully resident 32K point. Batching supplies BF16 with parallelism, while the 2.56× KV-byte reduction persists. The trend relative to FlashInfer holds under batching: at concurrency 8 the advantage is a 1.12× geometric mean across 8K–128K, from parity at 8K to 1.18–1.31× at 128K; for Qwen3-14B at concurrency 16 it is 1.09×/1.17× at the fully resident 16K/32K points. Points beyond 32K at concurrency 16 are capacity constrained because the KV pool admits only 11 and 5 requests at 64K and 128K, so we exclude them from aggregate results. Accuracy. Table 2 compares four backends across five benchmarks and two models. FP8 FlashInfer stays within 0.006 of BF16 for every model–task pair. Direct NVFP4 + FP8 loses most on Llama-3.1-8B GSM8K (0.108) and RULER-QA (0.091); MpFA recovers 62.5% of the loss on average (24%–83% per pair). It recovers 49%–77% on Llama-3.1-8B, trailing FP8 by at most 0.009 on knowledge and commonsense tasks. Qwen314B has smaller direct losses (0.006–0.015 on single-token tasks), with 50%–83% recovery on the four shorter tasks and 24% on RULER-QA. Across both models the residual concentrates on RULER-QA and GSM8K, where RULER-QA
10−2
0 16K
16K
32K
64K
sequence length
sequence length
(a) Relative overhead
(b) Absolute overhead
Figure 10. Cost of the two compensation placements for causal prefill with the Llama-3.1-8B shape. BF16 FA4 NVFP4+FP8
quantize / cache write MpFA(smoothing)
1.37×
30 25
1.22×
1.34×
5
1.08×
10
1.22×
15
1.00×
20
1.28×
1.00×
35
1.00×
prefill time per layer (ms)
40
0 16K
32K
64K
sequence length
Figure 11. Per-layer prefill cost, including attention and the quantization / cache-write pass (Llama-3.1-8B). recovery is 24%–49%: at 64K, error accumulated across many attention steps is less amenable to per-channel rank-one compensation. MpFA thus trades a few points on long-context and multi-step reasoning for a 2× lower-precision 𝑄𝐾 ⊤ path that is faster than FP8 FlashInfer (§7.2), while restoring most of the accuracy loss without retraining. 7.4
Ablation Study
Compensation placement. Figure 10 compares two placements of the same compensation term Δ𝑆 in QK4PV8 FA4. Staging compensation in shared memory, synchronizing after each KV block, and adding it element-wise incurs 20.1% 11
geometric-mean overhead across 16K–64K. With rank-one compensation, BF16 MMA first writes to tensor memory and NVFP4 𝑄𝐾 ⊤ accumulates into it, reducing overhead to 2.0%. At 64K, absolute overhead falls from 4.55 ms to 0.45 ms, or from more than half of the direct kernel’s prefill advantage to less than one tenth. Both placements compute the same term: element-wise placement uses FP32 scalar addition, whereas rank-one placement carries Δ𝑆 in BF16 and accumulates into FP32, adding one BF16 rounding with relative error no larger than 2−8 . Table 2 includes this effect because MpFA uses rank-one compensation by default. Quantization overhead. Quantization, sequence centering, and compensation generation remain part of end-to-end inference; MpFA fuses them with cache writes. Figure 11 includes this pass in the per-layer prefill cost and labels the resulting speedup over BF16 FA4. Across 32K and 64K, MpFA achieves a 1.25× geometric-mean speedup, compared with 1.34× for direct NVFP4 + FP8, preserving most of the kernel gain.
8
accuracy. These systems share a post-training, plug-andplay approach and do not require both matrix multiplications to use the same precision. SageAttention3 [30] extends this approach to consumer Blackwell GPUs, using NVFP4 matrix multiplication for both 𝑄𝐾 ⊤ and 𝑃𝑉 , together with 𝑄/𝐾 smoothing and probability quantization. HiFA4 [6] also provides training-free 4-bit attention, but targets the Ascend HIF4 NPU and its platform-specific formats and execution units. These systems do not address tensor memory, operand delivery, or asymmetric scaling between matmul and non-matmul units on data-center Blackwell GPUs. KVcache quantization methods such as KIVI [16] lower cache precision to save memory and traffic, complementing quantized attention that consumes a low-bit cache without dequantization. At the serving level, schedulers exploiting the distinct prefill and decode profiles—through phase disaggregation [21] or prefill-decode overlap [14]—are orthogonal to MpFA, which accelerates the attention operator each phase relies on. Different from prior work, MpFA is a high performance QK4PV8 attention kernel for data-center Blackwell GPUs. We characterize the B200 data path and hardware behavior of low-bit attention, build mixed-precision pipelines specialized for prefill and decode, and integrate them into an end-to-end serving system. MpFA further applies rank-one Tensor Core MMA smoothing compensation, which recovers accuracy without placing additional compensation work on the already saturated softmax path.
Related Work
FlashAttention [3, 4] reduces global-memory traffic through tiling and online softmax, and later work optimizes parallelism and data reuse for different GPUs. FlashAttention3 (FA3) [23] adds warp specialization on Hopper to overlap matrix multiplication and softmax across warps, and FlashAttention-4 (FA4) [27] redesigns the pipeline for Blackwell tensor memory, fifth-generation Tensor Cores, and asymmetric hardware scaling, serving as our direct B200 baseline. Other work adapts FlashAttention to the cache and datareuse characteristics of Arm multicore CPUs [9], while a parallel line designs dedicated attention accelerators, such as the approximation- and pruning-based sparse-attention architectures SpAtten [24] and ELSA [10]. Collectively, these systems show that attention performance depends on codesigning algorithmic tiling with the hardware data path rather than optimizing solely for matrix multiplication peak throughput. MpFA retains FlashAttention’s online softmax and FA4’s B200 execution model but addresses a different problem. While FA4 primarily uses BF16, MpFA moves 𝑄𝐾 ⊤ and 𝑃𝑉 to low-bit matrix multiplication paths and analyzes how block scaling, tensor memory capacity, and operand delivery jointly constrain the precision assignment. Prior low-bit LLM quantization methods, including MRGPTQ [7], AWQ [15], GPTQ [8], and SmoothQuant [25], target weights and activations in linear layers rather than attention activations. SageAttention [31] studies post-training quantization for attention: it smooths channel outliers in 𝐾, quantizes 𝑄𝐾 ⊤ to INT8, and keeps the more error-sensitive 𝑃𝑉 in FP16 with an FP16 accumulator. SageAttention2 [29] further quantizes 𝑄𝐾 ⊤ to INT4, represents 𝑃 and 𝑉 in FP8, and applies more extensive outlier smoothing to recover
9
Conclusion
We present MpFA, a training-free QK4PV8 FlashAttention kernel optimized for data-center Blackwell GPUs. Despite the B200’s substantially higher 4-bit Tensor Core throughput, the overheads of online quantization, data movement, and tensor memory contention make fully NVFP4 attention inefficient. MpFA instead adopts mixed-precision NVFP4(𝑄𝐾 ⊤ )+FP8(𝑃𝑉 ) and addresses its performance challenges through a finegrained asynchronous pipeline, tensor memory reuse, and stage-specific optimizations for prefill and decode. To reduce quantization error without further stressing the softmax path, MpFA introduces rank-one Tensor Core MMA smoothing compensation. On the B200, this compensation recovers 62.5% of the accuracy loss with only 2.0% average kernel overhead. Across 16K–128K contexts, MpFA achieves average prefill speedups of 1.31–1.33× over BF16 FA4 and 1.15–1.17× over FP8 FlashInfer for GQA 32:8 and 40:8, and achieves 2.81× the end-to-end output throughput of BF16 FA4 across both evaluated models.
References [1] Yushi Bai, Xin Lv, Jiajie Zhang, Hongchang Lyu, Jiankai Tang, Zhidian Huang, Zhengxiao Du, Xiao Liu, Aohan Zeng, Lei Hou, Yuxiao Dong, Jie Tang, and Juanzi Li. 2024. LongBench: A Bilingual, Multitask Benchmark for Long Context Understanding. In Proceedings of the 12
62nd Annual Meeting of the Association for Computational Linguistics (ACL). [2] Karl Cobbe, Vineet Kosaraju, Mohammad Bavarian, Mark Chen, Heewoo Jun, Lukasz Kaiser, Matthias Plappert, Jerry Tworek, Jacob Hilton, Reiichiro Nakano, Christopher Hesse, and John Schulman. 2021. Training Verifiers to Solve Math Word Problems. arXiv:2110.14168 (2021). [3] Tri Dao. 2024. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning. In International Conference on Learning Representations (ICLR). [4] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. 2022. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. In Advances in Neural Information Processing Systems (NeurIPS). [5] Chencheng Deng, Weiling Yang, Jianbin Fang, and Dezun Dong. 2026. Demystifying ARM SME to Optimize General Matrix Multiplications. In 2026 IEEE International Parallel and Distributed Processing Symposium (IPDPS). [6] Hui Dong, Yanzhao Li, Jie Gao, Chunlu Li, Zhiyuan Zhang, Yupeng Sun, Zhenyuan Chen, and Zhiqiang Zou. 2026. HiFA4: TrainingFree 4-bit FlashAttention on Ascend HIF4 NPUs for LLM Inference. arXiv:2607.04302 (2026). [7] Vage Egiazarian, Roberto L. Castro, Denis Kuznedelev, Andrei Panferov, Eldar Kurtic, Shubhra Pandit, Alexandre Marques, Mark Kurtz, Saleh Ashkboos, Torsten Hoefler, and Dan Alistarh. 2026. Bridging the Gap Between Promise and Performance for Microscaling FP4 Quantization. In International Conference on Learning Representations (ICLR). [8] Elias Frantar, Saleh Ashkboos, Torsten Hoefler, and Dan Alistarh. 2023. GPTQ: Accurate Post-Training Quantization for Generative Pre-trained Transformers. In International Conference on Learning Representations (ICLR). [9] Xiao Fu, Weiling Yang, Dezun Dong, and Xing Su. 2024. Optimizing Attention by Exploiting Data Reuse on ARM Multi-core CPUs. In Proceedings of the 38th ACM International Conference on Supercomputing (ICS). [10] Tae Jun Ham, Yejin Lee, Seong Hoon Seo, Soosung Kim, Hyunji Choi, Sung Jun Jung, and Jae W. Lee. 2021. ELSA: Hardware-Software Codesign for Efficient, Lightweight Self-Attention Mechanism in Neural Networks. In 2021 ACM/IEEE 48th Annual International Symposium on Computer Architecture (ISCA). 692–705. [11] Dan Hendrycks, Collin Burns, Steven Basart, Andy Zou, Mantas Mazeika, Dawn Song, and Jacob Steinhardt. 2021. Measuring Massive Multitask Language Understanding. In International Conference on Learning Representations (ICLR). [12] Cheng-Ping Hsieh, Simeng Sun, Samuel Kriman, Shantanu Acharya, Dima Rekesh, Fei Jia, Yang Zhang, and Boris Ginsburg. 2024. RULER: What’s the Real Context Size of Your Long-Context Language Models?. In Conference on Language Modeling (COLM). [13] Aaron Jarmusch and Sunita Chandrasekaran. 2026. Microbenchmarking NVIDIA’s Blackwell Architecture: An In-depth Architectural Analysis. arXiv:2512.02189 (2026). [14] Aditya K. Kamath, Ramya Prabhu, Jayashree Mohan, Simon Peter, Ramachandran Ramjee, and Ashish Panwar. 2025. POD-Attention: Unlocking Full Prefill-Decode Overlap for Faster LLM Inference. In Proceedings of the International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS). [15] Ji Lin, Jiaming Tang, Haotian Tang, Shang Yang, Wei-Ming Chen, Wei-Chen Wang, Guangxuan Xiao, Xingyu Dang, Chuang Gan, and Song Han. 2024. AWQ: Activation-aware Weight Quantization for OnDevice LLM Compression and Acceleration. In Proceedings of Machine Learning and Systems (MLSys). [16] Zirui Liu, Jiayi Yuan, Hongye Jin, Shaochen Zhong, Zhaozhuo Xu, Vladimir Braverman, Beidi Chen, and Xia Hu. 2024. KIVI: A TuningFree Asymmetric 2bit Quantization for KV Cache. In International Conference on Machine Learning (ICML).
[17] NVIDIA. 2022. NVIDIA H100 Tensor Core GPU Architecture. [18] NVIDIA. 2024. CUDA C++ Programming Guide, Version 12.4. [19] NVIDIA. 2024. NVIDIA Blackwell Architecture Technical Brief. Technical Report. NVIDIA. [20] Zaifeng Pan, Yitong Ding, Yue Guan, Zheng Wang, Zhongkai Yu, Xulong Tang, Yida Wang, and Yufei Ding. 2025. FastTree: Optimizing Attention Kernel and Runtime for Tree-Structured LLM Inference. In Proceedings of Machine Learning and Systems (MLSys). [21] Pratyush Patel, Esha Choukse, Chaojie Zhang, Aashaka Shah, Í nigo Goiri, Saeed Maleki, and Ricardo Bianchini. 2024. Splitwise: Efficient Generative LLM Inference Using Phase Splitting. In 51st Annual International Symposium on Computer Architecture (ISCA). [22] Keisuke Sakaguchi, Ronan Le Bras, Chandra Bhagavatula, and Yejin Choi. 2020. WinoGrande: An Adversarial Winograd Schema Challenge at Scale. In Proceedings of the AAAI Conference on Artificial Intelligence (AAAI). [23] Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, and Tri Dao. 2024. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision. In Advances in Neural Information Processing Systems (NeurIPS). [24] Hanrui Wang, Zhekai Zhang, and Song Han. 2021. SpAtten: Efficient Sparse Attention Architecture with Cascade Token and Head Pruning. In 2021 IEEE International Symposium on High-Performance Computer Architecture (HPCA). 97–110. [25] Guangxuan Xiao, Ji Lin, Mickael Seznec, Hao Wu, Julien Demouth, and Song Han. 2023. SmoothQuant: Accurate and Efficient Post-Training Quantization for Large Language Models. In International Conference on Machine Learning (ICML). [26] Zihao Ye, Lequn Chen, Ruihang Lai, Wuwei Lin, Yineng Zhang, Stephanie Wang, Tianqi Chen, Baris Kasikci, Vinod Grover, Arvind Krishnamurthy, and Luis Ceze. 2025. FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving. In Proceedings of Machine Learning and Systems (MLSys). [27] Ted Zadouri, Markus Hoehnerbach, Jay Shah, Timmy Liu, Vijay Thakkar, and Tri Dao. 2026. FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling. In Proceedings of Machine Learning and Systems (MLSys). [28] Rowan Zellers, Ari Holtzman, Yonatan Bisk, Ali Farhadi, and Yejin Choi. 2019. HellaSwag: Can a Machine Really Finish Your Sentence?. In Proceedings of the 57th Annual Meeting of the Association for Computational Linguistics (ACL). [29] Jintao Zhang, Haofeng Huang, Pengle Zhang, Jia Wei, Jun Zhu, and Jianfei Chen. 2025. SageAttention2: Efficient Attention with Thorough Outlier Smoothing and Per-Thread INT4 Quantization. In International Conference on Machine Learning (ICML). [30] Jintao Zhang, Jia Wei, Pengle Zhang, Xiaoming Xu, Haofeng Huang, Haoxu Wang, Kai Jiang, Jun Zhu, and Jianfei Chen. 2025. SageAttention3: Microscaling FP4 Attention for Inference and an Exploration of 8-Bit Training. In Advances in Neural Information Processing Systems (NeurIPS). [31] Jintao Zhang, Jia Wei, Pengle Zhang, Jun Zhu, and Jianfei Chen. 2025. SageAttention: Accurate 8-Bit Attention for Plug-and-Play Inference Acceleration. In International Conference on Learning Representations (ICLR). [32] Lianmin Zheng, Liangsheng Yin, Zhiqiang Xie, Chuyue Sun, Jeff Huang, Cody Hao Yu, Shiyi Cao, Christos Kozyrakis, Ion Stoica, Joseph E. Gonzalez, Clark Barrett, and Ying Sheng. 2024. SGLang: Efficient Execution of Structured Language Model Programs. In Advances in Neural Information Processing Systems (NeurIPS).
13