qsa_prompt_attn: tensor-core kernel for Volta (sm_70, mma.m8n8k4)

[fks]

On sm_70 qsa_prompt_attn_batch refuses the device (#371), so the prompt path runs the QSA attention on the decode kernel, one query at a time with FP32 FMAs: 19.2% of the GPU time of a 28,650-token prompt on a V100 (nsys, 0.1.35, attn_chunk_kernel<1>).

prompt_attn_v70_kernel is the v1 kernel's structure with the two matrix products on mma.sync.m8n8k4 (FP16 in, FP32 accumulate): q.k with one 64-dim group (= one int8 scale group) per warp and 8 cells per quad-pair, p.v with 16 dims per quad-pair. q and p keep their hi + lo FP16 split, and the four dim groups' partials are added in a fixed order (deterministic). Int8, FP16 and K8V4 KV are handled; Q4_0 KV (mode 4, sm_80+ only) keeps the old kernel. The m8n8k4 fragment layout was measured on a V100 and agrees with turbomind's SM70_MMA_884. It is used for compute capability 7.0 - 7.4; STRATA_PROMPT_ATTN_OLD=1 is the old path.

Accuracy (qsa_prompt_attn_parity, synthetic, V100, 32K context, 2,048 queries per chunk, output scale ~3.6):

KV	error vs FP64, old kernel	error vs FP64, new kernel	time per chunk, old -> new
int8	2.4e-6	5.1e-6	38.1 -> 13.4 ms (2.8x)
fp16	1.9e-6	5.3e-6	37.3 -> 28.3 ms (1.3x)

The 1,500- and 2,100-token contexts also pass. The test now skips its Q4_0 cases below sm_80 instead of exiting 2 there (the dispatcher keeps the old kernel for them).

Prompt speed (one V100-PCIE-32GB, i9-7960X, CUDA 12.8, origin/main at 99f3dbd, Unsloth UD-IQ4_XS, int8 KV; cold server, random-word prompts, greedy; same binary, STRATA_PROMPT_ATTN_OLD=1 against the default):

Prompt tokens	old path (tok/s)	new kernel (tok/s)
7,194	1,010	1,164	+15%
28,650	1,063	1,251	+18%
114,338	973	1,123	+15%

Decode is unaffected (it does not use this path).

Output checks

• Needle recall (tools/needle_bench.py, 8k / 32k / 128k at five depths, a fresh server per configuration): 15 of 15 found with the old path, 15 of 15 with the new kernel.
• Greedy output (48 tokens) is identical for the 28,650- and 114,338-token prompts. For the 7,194-token prompt it starts the same and differs later: another summation order in a greedy decode. I did not run a teacher-forced top-1 comparison.

Not covered: only the Unsloth IQ4_XS pack was run (not the GSQ-RCO quants); one request on one card. Method, profile and raw numbers: bench/results/2026-10-03-v100-prompt-attn/; build and kernel overview: docs/NVIDIA_V100.md (both in this PR). Independent of the BF16-via-FP16 change in #540: different code files, and it stacks with it (a user measured +9% on top of #540)

Written and measured with Claude Sonnet 5.5.

[sskver]

i tried it on top of #540 on a V100 (IQ2_XS, sm_70, CUDA 12.9), with only the kernel commit (ee0843d) cherry-picked onto my branch, which is upstream 0.1.37 plus the #540 commits

prefill, one run each:

• 8K prompt: 1104.8 -> 1205.5 tok/s (+9.1%), qsa attn 981 -> 524 ms
• 20K prompt: 1130.1 -> 1226.3 tok/s (+8.5%), qsa attn 2676 -> 1435 ms

so it stacks with #540. #540's attention commit changes qsa_decode_attn.cu, the kernel this PR bypasses during prefill, so the two don't overlap in code. that fits your +15-18% over the old path: #540's attention commit is worth about +7% on its own, and 1.07 x 1.09 is roughly 1.17

you mention you did not run a teacher-forced comparison, so for what it is worth, first-window logits against the same build without this kernel:

• 8K: argmax same, top-10 overlap 10/10 but order differs, KL 1.5e-3
• 20K: argmax same, top-10 overlap 9/10, order differs, KL 6.2e-3

that is the same size as the FP16 rounding in the BF16-via-FP16 change, so it doesn't look like a bug. if you have a KL or perplexity number on a longer prompt i'd like to see it, mine is one position per prompt and heavy-tailed

caveats: single runs, V100 only, i did not run ctest on this kernel

separately: the BF16 half of #540 (BF16 GEMMs through FP16 tensor cores, same STRATA_BF16_VIA_F16 toggle) is the same change as #593, so i left a note there. this PR itself doesn't overlap with either

[fks]

Thanks for testing this. I measured the teacher-forced comparison on longer prompts, with the project's method (STRATA_LOGPOS, top-256).

Setup: a chat whose first message (8.5k / 33k / 99k tokens of this repository's docs and code) is read by the batched prompt path, then a ~600-token last message that goes through the verify windows and is scored (617 positions per prompt). One V100, UD-IQ4_XS, old path = STRATA_PROMPT_ATTN_OLD=1. KL is old || other over the old run's top 256 plus a bucket for the rest.

comparison	KL mean (8.5k / 33k / 99k)	argmax same	top-10 overlap
old vs old (rerun), new vs new (rerun)	0	100%	100%
old vs new kernel	0.028 / 0.040 / 0.116	97.7 / 93.2 / 86.4%	93.3 / 91.9 / 88.4%
control: BF16 GEMMs without the FP16 detour	0.026 / 0.051 / 0.090	97.2 / 91.4 / 91.2%	92.9 / 91.4 / 88.5%
control: --prefill 2048 instead of auto	0.023 / 0.035 / 0.038	97.9 / 94.8 / 95.6%	92.9 / 92.2 / 89.7%

So the model amplifies small numeric differences a lot (heavy-tailed: the median KL at 8.5k is 7e-4, p99 0.35), and both controls, which change nothing but the rounding order, move the distribution as much as the kernel does. At 99k the kernel is somewhat above both controls; that is one document and 617 positions, so I can neither confirm nor exclude a real effect there.

Directly on the model's data, with a temporary check that ran the decode kernel beside the new one on the same inputs (every 8th query batch, all 12 QSA layers, ~5 M (query, head) rows, positions up to 98k): ||new - old|| / ||old|| per row is 4.4e-7 on average, 4.8e-6 in the worst row, max absolute difference 4e-5 against values up to 31, with no trend over layers or context length. That is FP32 rounding level.

Caveats: V100 only, one run per setting (the reruns are bit-identical), the check code is not part of the PR. The table and the method are now in bench/results/2026-10-03-v100-prompt-attn/README.md (dbb1a08).

[sskver]

that row-level check pretty much settles it for me, only thing i noticed is 99k sitting a bit above the controls. looks good to me!

[Niko1221]

Merged for the next release, with your commits and docs/NVIDIA_V100.md. The kernel and its launcher are compiled into the experimental build only (-DSTRATA_EXPERIMENTAL_SM60=ON), so the ready-made engine loads no extra kernels. Thanks for the teacher-forced comparison with controls.