Long-context modeling is a pivotal capability for Large Language Models, yet the quadratic complexity of attention remains a critical bottleneck, particularly during the compute-intensive prefilling phase. Our previous work, FlashPrefill, mitigates this cost through instantaneous pattern discovery and max-based dynamic thresholding; however, it remains an algorithmic prototype that is still distant from production deployment. In this paper, we present FlashPrefill V2, which evolves FlashPrefill from a prototype toward practical long-context serving along three dimensions. First, we introduce a mean correction term that effectively suppresses the approximation error, keeping performance degradation manageable even at extreme sparsity levels. Second, we redesign the sparse attention operator with PackGQA memory access, warp specialization, and pingpong pipelining, fully aligning with the latest FlashAttention-3/4 implementations and supporting FP8 inference to meet practical quantization requirements. Third, FlashPrefill V2 natively supports paged KV cache and continuous batching, allowing integration as an attention backend in modern inference frameworks such as SGLang. Extensive evaluations on NVIDIA H20 GPUs---among the most widely deployed inference accelerators---demonstrate that FlashPrefill V2 delivers up to 47.26x and 27.19x speedups over FlashAttention-2 at 128K context length under FP8 and BF16 precision, respectively, and, in FP8, still achieves a 30.49x speedup against an FA3/4-aligned dense baseline.
Figures & tables
Figure 1: FlashPrefill V2 evolves FlashPrefill [ 8 ] from an algorithmic prototype toward practical long-context serving along three dimensions: (1) a mean correction term that preserves model accuracy under extreme sparsity; (2) an FA3/4-aligned sparse attention kernel with PackGQA memory access, warp specialization, pingpong pipelining, and FP8 support; and (3) native compatibility with paged KV cache and continuous batching for integration as an attention backend in modern serving frameworks such as SGLang [ 42 ] .
Figure 2: Speedup of various attention operators relative to FlashAttention-2 [ 5 ] on NVIDIA H20 GPUs, including an FA3/4-aligned dense baseline [ 22 , 38 ] . All results are measured with a batch size of 4. FlashPrefill V2 exhibits a dominant advantage, particularly in long-context scenarios, and its FP8 variant further amplifies the gains.
Figure 3: The FlashPrefill V2 prefill pipeline. Stage 1 packs queries with PackGQA, pools block-mean statistics, and produces a CSR index of selected blocks in a single fused pass; Stage 2 runs the warp-specialized sparse attention kernel over the selected blocks (exact path) while the unselected blocks are compensated through the mean correction path (Sec. 3.2 ), and the two streams merge inside the online softmax.
Figure 4: Mean correction. Selected blocks (red) are computed exactly, while each pruned block is pooled into its mean statistics (kˉJ,vˉJ) and contributes a surrogate term ∣BJ∣esˉJ(vˉJ,1) to the softmax numerator and denominator, recovering the discarded probability mass without per-token computation.
Method
RULER Score
Operator Speedup over FA2
4K
8K
16K
32K
64K
128K
Avg.
4K
8K
16K
32K
64K
128K
Llama-3.1-8B-Instruct
Full
95.98
94.74
93.07
89.06
86.27
73.82
88.82
1.00 ×
1.00 ×
1.00 ×
1.00 ×
1.00 ×
1.00 ×
MInference
95.06
94.23
92.18
87.68
83.01
71.61
87.30
0.12 ×
0.18 ×
0.46 ×
0.83 ×
1.34 ×
2.45 ×
FlexPrefill
95.18
94.09
92.35
86.79
84.08
72.01
87.42
0.12 ×
0.33 ×
0.98 ×
2.21 ×
4.16 ×
5.18 ×
XAttention
95.24
94.37
91.61
87.03
83.65
72.33
87.37
0.79 ×
1.29 ×
1.83 ×
2.34 ×
3.19 ×
3.48 ×
Table 1: RULER scores (left) and attention operator speedups over FlashAttention-2 (right) at sequence lengths from 4K to 128K. The speedups are measured end-to-end on needle-in-a-haystack inputs, including pattern discovery, thresholding, and mean correction. Results marked with ∗ (in gray) are measured with direct online FP8 quantization without corrected weights.
Method
Single-Document QA
Multi-Document QA
Summarization
Few-shot Learning
Synthetic
Code
Avg.
NarrQA
Qasper
MF-en
MF-zh
Hotpot
2Wiki
MuSiQue
DuRead
GovRep
QMSum
MNews
VCSum
TREC
TQA
SAMSum
LSHT
Count
PR-en
PR-zh
LCC
RepB
Llama-3.1-8B-Instruct
Full
29.39
44.61
56.34
63.63
57.48
48.08
32.28
34.43
34.51
25.29
26.80
17.37
72.50
91.48
43.59
46.50
11.10
100.00
90.45
62.93
56.19
49.76
MInference
28.13
42.36
53.67
62.08
56.28
46.72
32.01
33.41
33.28
25.16
27.02
16.84
69.50
90.83
42.87
43.00
6.00
96.00
84.00
60.41
55.92
47.88
FlexPrefill
27.62
42.67
52.98
62.34
57.14
45.32
30.97
33.72
33.87
24.58
26.34
17.02
67.50
90.36
44.72
43.00
6.00
94.00
83.33
61.72
55.83
47.67
XAttention
28.52
43.12
54.01
60.19
56.73
45.98
32.17
33.15
34.55
25.37
27.21
17.41
71.00
91.37
44.05
44.00
6.88
95.50
85.68
61.05
55.36
48.25
Table 2: Performance comparison on all 21 tasks of LongBench [ 2 ] , grouped by task category. Results marked with ∗ (in gray) are measured with direct online FP8 quantization without corrected weights.
Model
4K
8K
16K
32K
64K
128K
Llama-3.1-8B
76.0%
54.2%
33.6%
18.7%
9.0%
4.6%
Qwen3-4B
72.0%
48.5%
29.8%
17.6%
9.2%
4.9%
Qwen3-30B-A3B
70.4%
46.0%
29.6%
16.2%
9.4%
4.9%
Table 3: Attention density of FlashPrefill V2 measured on needle-in-a-haystack inputs at each sequence length.
BSZ
FA3/4
FlashPrefill V2
FlashPrefill V2-FP8
4K
8K
16K
32K
64K
128K
4K
8K
16K
32K
64K
128K
4K
8K
16K
32K
64K
128K
Llama-3.1-8B-Instruct
1
0.15
0.29
0.63
1.49
3.95
11.81
0.15 (1.02 × )
0.29 (1.02 × )
0.58 (1.09 × )
1.17 (1.28 × )
2.46 (1.60 × )
5.49 (2.15 × )
0.10 (1.53 × )
0.18 (1.63 × )
0.35 (1.81 × )
0.71 (2.09 × )
1.48 (2.66 × )
3.23 (3.66 × )
4
0.58
1.17
2.54
5.94
12.30
32.94
0.57 (1.01 × )
1.14 (1.03 × )
2.32 (1.09 × )
4.79 (1.24 × )
7.69 (1.60 × )
15.39 (2.14 × )
0.35 (1.66 × )
0.68 (1.73 × )
1.36 (1.87 × )
2.80 (2.12 × )
4.49 (2.74 × )
8.84 (3.73 × )
8
1.07
2.21
4.87
9.83
21.23
56.67
1.06 (1.01 × )
2.14 (1.03 × )
4.39 (1.11 × )
7.77 (1.27 × )
13.18 (1.61 × )
26.13 (2.17 × )
0.67 (1.58 × )
1.39 (1.59 × )
2.81 (1.73 × )
4.51 (2.18 × )
7.66 (2.77 × )
14.93 (3.80 × )
16
2.60
3.85
8.14
16.10
37.12
102.84
2.58 (1.01 × )
3.59 (1.07 × )
7.31 (1.11 × )
12.56 (1.28 × )
22.93 (1.62 × )
46.96 (2.19 × )
1.44 (1.80 × )
2.97 (1.30 × )
4.62 (1.76 × )
7.28 (2.21 × )
13.25 (2.80 × )
26.82 (3.83 × )
Table 4: End-to-end time-to-first-token (TTFT, seconds) of Llama-3.1-8B-Instruct, Qwen3-4B-Instruct-2507, and Qwen3-30B-A3B-Instruct-2507 served by SGLang on the needle-in-a-haystack workload. Each entry is averaged over 30 repeated measurements. Parenthesized numbers are speedups over the FA3/4 backend.
Rate (req/s)
Backend
TTFT (s)
TPOT (ms)
Throughput
P50
P99
P50
P99
(tok/s)
(req/s)
Qwen3-4B-Instruct-2507
1
FA3/4
76.73
154.64
419
1170
24.0
0.37
FlashPrefill V2
17.23 (4.45x)
48.73 (3.17x)
258 (1.62x)
601 (1.94x)
45.5 (1.90x)
0.71
FlashPrefill V2-FP8
2.79 (27.47x)
17.27 (8.95x)
47 (8.84x)
487 (2.40x)
58.4 (2.44x)
0.91
4
FA3/4
79.05
161.54
513
1102
23.8
0.37
Table 5: Open-loop serving results on Qwen3-4B-Instruct-2507 and Qwen3-30B-A3B-Instruct-2507 under Poisson arrivals with mixed prompt lengths (4K–128K RULER needle-in-a-haystack, 100 requests per cell, 64 output tokens each). Parenthesized numbers are speedups over the FA3/4 backend.
Chunk size
Backend
TTFT (s)
TPOT (ms)
Throughput
P50
P99
P50
P99
(tok/s)
(req/s)
Qwen3-4B-Instruct-2507
8K
FA3/4
80.98
165.14
597
1228
22.5
0.35
FlashPrefill V2
50.15 (1.61x)
95.23 (1.73x)
360 (1.66x)
684 (1.80x)
38.7 (1.72x)
0.60
FlashPrefill V2-FP8
33.47 (2.42x)
61.20 (2.70x)
259 (2.30x)
529 (2.32x)
60.7 (2.70x)
0.95
16K
FA3/4
79.42
161.96
580
1203
22.9
0.36
Table 6: Open-loop serving at 16 req/s with chunked prefill at chunk sizes of 8 K and 16 K (same workload as Tab. 5 , where chunking is disabled). Parenthesized numbers are speedups over the FA3/4 backend at the same chunk size.
Figure 11
HPC-Ops BSA
FlashPrefill V2
Sparse index format
dense mask, padded, per Q head
CSR of selected blocks, per KV head
Index memory @ 64K / 128K
8 MB / 32 MB (any density)
3 MB / 6 MB (9% / 5% density)
GEMM–softmax overlap
serial
intra-warpgroup pingpong
Paged KV addressing
per-page TMA (page size ∣ 128)
cp.async, arbitrary page size
KV splitting
none
count-balanced CSR splits
Mean correction
none
fused into the mainloop
Table 7: Design comparison with the HPC-Ops block-sparse attention kernel. Index memory is evaluated at batch size 1 with 32 query heads and 4 KV heads: the dense mask stores one byte per (query head, 128 -token Q tile, 128 -token KV tile) regardless of density, whereas the CSR index stores only the selected 64 -token tiles as int32 at the measured needle-in-a-haystack densities (Tab. 3 ), is organized per KV head under PackGQA, and matches the SGLang index format. The fused correction path has no counterpart in HPC-Ops BSA.
Method
4K
8K
16K
32K
64K
128K
Avg.
Full Attention
94.82
92.98
91.68
88.86
82.81
71.22
87.06
V2 w/o correction
94.38
92.06
91.18
87.21
80.92
68.86
85.77
FlashPrefill V2
94.49
92.42
91.56
87.48
81.67
69.76
86.23
V2-FP8 w/o correction
93.78
91.82
90.68
85.04
76.86
62.78
83.49
FlashPrefill V2-FP8
94.13
92.37
91.36
87.47
80.59
68.97
85.82
Table 8: Ablation of the mean correction term on RULER with Qwen3-4B-Instruct-2507, in both BF16 and FP8. Discarding unselected blocks without compensation degrades accuracy increasingly toward longer contexts, while the zero-order correction stays close to full attention at all lengths in both precisions.
α
0.2
0.1
0.05
0.025
0.0125
Density (%)
5.2
9.2
14.1
19.8
23.6
V2-FP8 w/o correction
74.68 ( − 5.44)
76.86 ( − 3.73)
77.16 ( − 3.51)
77.41 ( − 3.31)
78.02 ( − 2.73)
FlashPrefill V2-FP8
80.12
80.59
80.67
80.72
80.75
Table 9: Ablation of the selection threshold α on RULER at 64K with Qwen3-4B-Instruct-2507 in FP8. Density is the fraction of selected attention blocks. Red numbers are the score drops of the uncorrected variant relative to the corrected pipeline at the same α . Full attention scores 82.81 .
Long-context modeling is a pivotal capability for Large Language Models, yet the quadratic complexity of attention remains a critical bottleneck, particularly during the compute-intensive prefilling phase. While various sparse attention mechanisms have been explored, they typically suffer from either significant search latency or insufficient sparsity. In this paper, we propose FlashPrefill, a framework enabling ultra-fast prefilling via instantaneous pattern discovery and thresholding. FlashPrefill leverages a fast block-searching technique to simultaneously locate dynamic vertical, slash, and block-sparse attention patterns. Crucially, it introduces a dynamic thresholding mechanism that bypasses the prohibitive overhead of sorting or accumulating attention scores while effectively eliminating the long-tail distribution to enhance sparsity. Extensive evaluations demonstrate that FlashPrefill achieves a substantial leap in efficiency, delivering an unprecedented 27.78x speedup on 256K sequences. Notably, unlike existing methods that incur efficiency degradation on shorter contexts, FlashPrefill maintains a 1.71x speedup even at a 4K context length, demonstrating its robustness and practical utility across varying sequence scales.
The quadratic complexity of self-attention during the prefill phase impedes long-context inference in large language models. Existing sparse attention methods face a trade-off among context adaptivity, sampling overhead, and fine-tuning costs. We propose VSPrefill, a mechanism requiring lightweight training that uses the vertical-slash structural pattern in attention distributions. Our compact VSIndexer module predicts context-aware importance scores for vertical columns and slash diagonals from key-value representations augmented with RoPE. This approach constructs sparse masks with linear complexity without modifying the backbone parameters. During inference, an adaptive cumulative-threshold strategy allocates sparsity budgets per layer, while a fused kernel executes attention with on-the-fly index merging. Evaluated on Qwen3-4B-Instruct and LLaMA-3.1-8B-Instruct across the LongBench and RULER benchmarks, VSPrefill preserves 98.35% of the full attention accuracy while delivering a 4.95x average speedup at a context length of 128k. These results establish a new Pareto frontier in the trade-off between accuracy and efficiency.
Chen Guanzhong
School of Computer Science and Technology, Tongji University
Modern large language models increasingly require long contexts for reasoning and multi-document tasks, but attention's quadratic complexity creates a severe computational bottleneck. We present Block Sparse Flash Attention (BSFA), a drop-in replacement that accelerates long-context inference while preserving model quality. Unlike methods that predict importance before computing scores, BSFA computes exact query-key similarities to select the top-k most important value blocks for each query. By comparing per-block maximum scores against calibrated thresholds, we skip approximately 50% of the computation and memory transfers for pruned blocks. Our training-free approach requires only a one-time threshold calibration on a small dataset to learn the per-layer and per-head attention score distributions. We provide a CUDA kernel implementation that can be used as a drop-in replacement for FlashAttention. On Llama-3.1-8B, BSFA achieves up to 1.13x end-to-end speedup on LongBench with only a 1.1% accuracy drop, and up to 1.24x on Needle-in-a-Haystack retrieval at a 1% accuracy drop. The attention kernel itself accelerates by up to 1.38x. We compare BSFA against five recent sparse attention baselines (SpargeAttention, MInference, FlexPrefill, XAttention, and BLASST), and verify the method on Qwen2.5-7B and on A6000 and H100 GPUs. The implementation is available at https://github.com/Danielohayon/Block-Sparse-Flash-Attention.
Daniel Ohayon, Itay Lamprecht, Itay Hubara +3
Technion – Israel Institute of Technology, Haifa, Israel · Intel - Habana Labs · Intel – Habana Labs, Tel Aviv, Israel