Repository navigation
sm_75: the v2 int8 prompt attention with register staging instead of cp.async (STRATA_PA_V2T=1) - #2092
Open
rafatxf wants to merge 1 commit into
Open
sm_75: the v2 int8 prompt attention with register staging instead of cp.async (STRATA_PA_V2T=1)#2092rafatxf wants to merge 1 commit into
rafatxf wants to merge 1 commit into
Conversation
…cp.async (STRATA_PA_V2T=1) On Turing the int8-KV prompt attention ran v1 because v2 stages its K/V slices with cp.async. Here a warp loads its next slice into registers and stores it into its other stage after the current chunk (the same bytes in the same places); the MMAs are the m16n8k8 pairs v1 already uses. Accuracy as v1 (qsa_prompt_attn_parity vs FP64: 3.7e-6 / 4.9e-6 / 4.5e-6 / 3.8e-6 of 3.5-3.7). RTX 2080 Ti, ms per chunk v1 -> v2: 32K x 1024 queries 13.0 -> 7.6, 64K x 512 6.7 -> 4.8, 1.5K 6.2 -> 5.0, 2.1K x 256 3.2 -> 2.3. Opt-in until the engine A/B. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
This was referenced Oct 12, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
On Turing the int8-KV prompt attention runs v1, because v2 (
prompt_attn_i8_kernel) stages each warp's K/V slices withcp.async, which does not exist before sm_80. v1 keeps one 4-warp CTA per SM (37 KB of shared memory) and gathers a chunk, waits, computes; the attention is ~18 % of a 130K prompt read on an RTX 2080 Ti.Here, below sm_80, v2's
issue()loads the warp's next pieces into registers (4 + 4uint4a lane) andstore_stage()writes them into the warp's other stage after the current chunk's p·v, where v2 already stores the next chunk's scales - the same bytes in the same places as the cp.async pipeline; each warp only touches its own stages, so__syncwarpsuffices as before. The MMAs are the m16n8k8 pairs v1 already uses on Turing. Opt-in (STRATA_PA_V2T=1): another summation order than v1, like v2 against v1 on sm_80.Setup
RTX 2080 Ti 22 GB (sm_75), i7-6700, CUDA 13.0, driver 595, SM 1650 MHz, 240 W;
v0.1.43. Model Q2_0,--resident-experts,--max-context 524288,--kv int8 --kv-resident 32768,STRATA_BF16_TC=1,STRATA_SELECT_SIMT=2.qsa_prompt_attn_parity(int8 KV), ms per chunk and error against FP64Engine
98K prompt read (one-shot, two runs each): 1473 → 1550 tok/s (+5.2 %), deterministic run to run.
Quality, teacher-forced as in #655 / #2071 (
--serve, the first P tokens through the prompt path, the next 2,048 resumed through the verify windows,STRATA_LOGPOS_TOPK=20, raw Spanish text):Noise floor: at 32K two v1 runs are identical; at 98K two v1 runs already differ by KL 0.0103 (argmax 95.5 %), so most of the 98K difference is the engine's own run-to-run variation there. Perplexity moves by −1.8 % and −0.1 %.
The server here runs with
STRATA_PA_V2T=1since 2026-10-12.🤖 Generated with Claude Code