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

Beschreibung

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.