#2197 Turing (sm_75, RTX 20) support: tracking issue for the open PRs

Descripción

Over the last days I opened nine PRs that make Strata run better on Turing (RTX 2080 Ti here). Each one stands alone with its own measurements, but together they are easier to review as a set, so this issue lists them, how they depend on each other, and what they add up to. Everything was measured on one card; I have no other sm_75 card to check on.

## Setup

RTX 2080 Ti 22 GB (sm_75), i7-6700 (4c/8t, AVX2), 56 GB DDR4-2133, CUDA 13.0, SM locked at 1650 MHz, 240 W. Base `v0.1.43` (7695859a). Model Qwen3.8-Flash-Next, canonical Q2_0 pack, `--max-context 524288`, `--prefill auto:16384`, served with `--serve`.

## The PRs

| PR | area | what | bitwise vs today | on by default? |
|---|---|---|---|---|
| #2069 | decode | Turing tables for the decode matvecs (`kIlRows_75`) and the grouped experts (mode 4) | yes | Turing default, switch to turn off |
| #2070 | decode | hyper-connection up kernel for 3+ tokens, staged down kernel in token pairs | yes | Turing default, switch |
| #2072 | decode | decode QSA selection at long context: halving scorer + split top-k | yes | Turing default, switch |
| #2071 | prompt | tiled FP32 block scorer `STRATA_SELECT_SIMT=2` (+ the case for SIMT on RTX 20 by default) | same bits as `=1` | opt-in |
| #2093 | prompt | prompt top-k with one full read of the scores (identical ids) | yes | Turing default — **stacked on #2072** |
| #2091 | prompt | hyper-connection `xn16` written as the FP16 its tensor-core GEMMs read | yes | with `STRATA_BF16_TC=1` |
| #2092 | prompt | v2 int8 prompt attention with register staging instead of `cp.async` | no (FP32-level, KL-checked) | opt-in `STRATA_PA_V2T=1` |
| #2146 | prompt | that v2 attention in one shared-memory stage: 3 blocks per SM instead of 1 | yes (vs #2092) | follows #2092 — **stacked on #2092** |
| #2145 | prompt | the fused int8 prompt experts (`moe_fused.cu`, `moe_fused_iq.cu`) ported to sm_75, Q2_0 gate/up | no (int8 rounding, as on sm_80; KL-checked) | opt-in `STRATA_PF_FUSED=1 STRATA_PF_FUSED_SM75=1` |

Suggested order: #2069, #2070, #2072 → #2093; #2071; #2091; #2092 → #2146; #2145. They touch different files except where noted as stacked.

## Together

All of them on (`STRATA_BF16_TC=1 STRATA_SELECT_SIMT=2 STRATA_PA_V2T=1 STRATA_PF_FUSED=1 STRATA_PF_FUSED_SM75=1`), against `v0.1.43` with its defaults:

| | v0.1.43 | with the PRs |
|---|---:|---:|
| 389,560-token prompt from scratch | 517.6 s | 223.4 s |
| +16K tokens appended at 406K | 31.5 s | 11.4 s |
| decode, 4K context | 62.8 t/s | 66.7 t/s\* |
| decode after the 389,560-token prompt | 52.8 t/s | 62.6 t/s\* |
| needle (2 facts) at 352,660 tokens | — | 2/2, 199 s |

\* `--serve`, 512 new tokens after the prompt, two runs each (MTP drafting, `--spec 4 --spec-min-p 0.7`; the rate moves with the text by ~±2 %). The v0.1.43 column is its own binary with no environment set; the prompt times are the same sessions'.

Quality was checked per PR (bitwise ones by hashes and first logits; the two that are not by teacher-forced KL / perplexity against the existing path and against the FP16 expert path — numbers in #2092 and #2145).

## Also measured here (settings only, no code change)

With the PRs above in place, three existing options moved decode further on this card (512-token generations, two runs each):

- `--adapt-async 1 --adapt-every 2`: in the resident RAM mode the blocking adaptive tier stopped the GPU for ~7 ms about every third window (nsys: 10 % of decode time idle, the stalls holding the 1.38 MB blob copies). Async: +3.6 % at 4K, +7.4 % at 128K, +2.9 % at 390K; a round every window was no better. (`STRATA_ADAPT_OVERLAP_RESIDENT=1` was −4.5 % here, the copies competing with the CPU experts for DDR4-2133.)
- `STRATA_Q2_BITPLANE=1`: +2 to +4 % once the tier no longer stalls (the CPU experts are then on the critical path of an i7-6700).
- A draft vocabulary built with `tools/draft_vocab.py --corpus` from the model's own answers in Spanish (English/code base + 2,735 tokens, 43,260 ids, 73 MiB instead of the 106K-id default's 184 MiB): 21 % of the tokens in those answers — Markdown included — were outside the English/code base. +7 % in Spanish, +4 % on code, +7 % on English text at 32K.

## A question for you

If these look right, would you consider making the opt-in ones (`STRATA_SELECT_SIMT=2`, `STRATA_PA_V2T=1`, `STRATA_PF_FUSED_SM75=1`, and `STRATA_BF16_TC=1` from #655) the defaults on sm_75? Happy to change any of them to fit how you prefer the code, or to measure anything else on this card.

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