#2092 sm_75: the v2 int8 prompt attention with register staging instead of cp.async (STRATA_PA_V2T=1)

Descrição

On Turing the int8-KV prompt attention runs v1, because v2 (`prompt_attn_i8_kernel`) stages each warp's K/V slices with `cp.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 + 4 `uint4` a lane) and `store_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 `__syncwarp` suffices 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 FP64

| context × queries | v1 | v2 (register staged) | error vs FP64, v1 / v2 (output scale) |
|---|---:|---:|---|
| 32K × 1024 | 12.98 | 7.63 | 4.3e-6 / 3.7e-6 (3.7) |
| 64K × 512 | 6.73 | 4.79 | 3.7e-6 / 4.9e-6 (3.5) |
| 1.5K × 1500 | 6.20 | 5.03 | 4.5e-6 / 4.5e-6 (3.6) |
| 2.1K × 256 | 3.16 | 2.31 | 3.3e-6 / 3.8e-6 (3.6) |

## Engine

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):

| prefix | argmax agreement | KL (v1 ‖ v2) | perplexity v1 / v2 | argmax where v1's top-2 gap ≥ 0.5 |
|---|---:|---:|---:|---:|
| 32,768 | 97.0 % | 0.0169 | 4.623 / 4.538 | 1937 / 1963 |
| 98,304 | 94.3 % | 0.0137 | 5.426 / 5.421 | 1661 / 1675 |

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=1` since 2026-10-12.

🤖 Generated with [Claude Code](https://claude.com/claude-code)