#627 opt for sm_70(just test V100-16g*1 or 2)

Descripción

This branch adds a set of runtime optimizations for the Tesla V100 (compute capability 7.0) and puts all of them behind one CMake switch, `STRATA_EXPERIMENTAL_V100`, in the same style as the existing `STRATA_EXPERIMENTAL_SM60`. The switch defaults to OFF: the optimization sources are not compiled, every entry point takes the original path, and trunk behavior (including numerics) stays bit-for-bit unchanged. With the switch ON all optimizations are active, and `-DSTRATA_EXPERIMENTAL_V100=ON` alone gives a runnable optimized sm_70 build (it also admits sm_70, equivalent to enabling `STRATA_EXPERIMENTAL_SM60` alongside). Each optimization keeps a fine-grained runtime A/B knob (`STRATA_QPN8`, `STRATA_SPLIT_P2P`, `STRATA_GROUPED_ATTN`, `STRATA_GDN_CHUNK`, `STRATA_PREFILL_BF16_F16`, ...) so any one of them can be verified alone.

## How to build

```bash
cmake -DSTRATA_ENABLE_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=70 -DSTRATA_EXPERIMENTAL_V100=ON ...
```

or through the installer (mirroring the #295 `STRATA_EXPERIMENTAL_SM60=1` flow):

```bash
STRATA_EXPERIMENTAL_V100=1 python3 setup.py --build ...
```

A CUDA 12.x toolkit is required (CUDA 13 dropped offline compilation for sm_70).

## Optimizations and measurements

| Optimization | Where | Measured (V100-SXM2-16GB) | Numerics |
| --- | --- | --- | --- |
| Volta tensor-core port of the QSA prompt attention (mma.sync m8n8k4 / wmma m16n16k16, hi/lo FP16 split; no cp.async, no ldmatrix, row padding instead of XOR-swizzle) | `qsa_prompt_attn_sm70.cu` (new), dispatch gated to 70≤cc<75 | 2.2× at the full 32K prompt shape (int8 KV 34.4→15.8 ms) | FP32-level (same contract as qsa_prompt_attn.hpp), new-vs-old ~2e-6 of scale |
| prefill BF16 GEMMs go through FP16 on cc<8.0 (V100 has no BF16 ALUs; exact bf16→f16 conversion + W twin cache + a scratch-tiling fix) | `gemm.cu`, `bf16_bits.hpp` | GEMM itself 4.3–21.7×; end-to-end prefill 1615→1816 tok/s | conversion exact; GEMM FP32-level |
| operator fusions for decode GEMV / MoE (gate/up+SwiGLU in one kernel, multi-weight GEMVs, activation-quantize epilogue) | `s_gemv_pair.cu` (new), `quantize_act.cu`, `bf16_gemv.cu`, `layer.cpp` | kernel time -26% to -51% | **bitwise** (memcmp clean) |
| grouped attention for MTP verify windows (one launch per M=2..8 group, Volta WMMA, each KV row read once) | `qsa_grouped_attn.cu` (new) | 1.14×–3.37× per window (2.7–3.4× vs row-by-row) | FP32-level, ≤1e-3× output scale |
| layer-split hand-off over NVLink P2P (was a pinned-RAM bounce; peer store inside the graph for verify, cudaMemcpyPeerAsync for prompt chunks) | `device.cu`, the hand-off buffers in `verify.cpp`/`prefill.cpp` | hand-off bandwidth 3.3→48 GB/s (nv2 link, measured) | **bitwise** (A/B identical) |
| vectorized int8 KV gather (8 int8 per thread, uint2/uint4 loads and stores) | `kv_q8.cu` | 32K/2051-cell gather 401→452 GB/s (+11.4%) | **bitwise** |
| GDN recurrence chain-split (default on sm_70; the chunked FlashQLA path stays opt-in) | `prefill/kernels.cu`, `gdn_chunk.cu` | recurrence kernel +6–7% | FP32-level (bit-exact variant kept) |
| QPN8 expert GEMV repack (m8n8k4 tensor cores, load-time weight repack + per-window A-fragment prep; pays off only at M≥4, opt-in `STRATA_QPN8=1`) | `s2_qpn8.cu` (new), `expert_cache.cpp` | full pipeline M=8 1.19–1.22× | quantized intermediates byte-identical, FP32 summation order bound |

Two bugs fixed along the way (always on, not gated): `mtp.cpp` allocated one row stride too few for the batch attention scratch (overrun risk); the grouped attention cached its shared-memory opt-in per process, so under `--layer-split` device1 reused device0's result and graph capture failed.

## 32k-context benchmarks

Model Swift 1.5 Qwen3.8-Flash-Next GSQ-RCO IQ2_XS (the two shards in ~/Qwen/ukisai), context 32768 (32504 prompt tokens), int8 KV, `--prefill auto`, `--spec 4 --spec-min-p 0.5` with the MTP draft layer, greedy, `--expert-cache auto`; the two-card runs use `--layer-split auto`. Hardware: 2× Tesla V100-SXM2-16GB (NVLink nv2), AMD EPYC 7F52, 114 GB RAM, CUDA 12.8, driver 580.173.02. Protocol: one warm-up round per engine, then 5 rounds, median reported; prefill counts fresh tokens (prompt cache off); decode is the median of the rounds that ran the full 256 tokens.

| Configuration | prefill tok/s | decode tok/s (median of full 256-token rounds) |
| --- | ---: | ---: |
| v100-opt single card (this branch, switch ON) | **1473** | 32.0 |
| v100-opt two cards (this branch, switch ON) | **2087** | 49.8 |
| main single card (reference; same as switch OFF) | 1115 | 34.1 |
| main two cards (reference; same as switch OFF) | 1615 | 52.4 |

prefill vs main: +32.7% on one card, +29.2% on two; two cards vs one card (this branch) +41.7%. decode between the two builds falls inside the measurement noise (see below).

A note on the decode numbers: greedy generation is not round-deterministic in this configuration — expert-cache admission is timing-dependent (the CPU pool and the GPU cache round the same expert differently; bench/results/2026-09-29-layer-split documents the effect), so rounds generate different text, the MTP draft acceptance rate swings between 58% and 94%, and end-to-end decode tok/s varies about ±15% between sessions. The table therefore reports full-length rounds only, and the single-card figures rest on few such rounds (most rounds hit EOS early); treat them as magnitude references. The per-phase kernel gains (measured column above) are solid, but at 32k the end-to-end decode is bounded by expert bandwidth and the CPU pool, so the improvement is hidden inside that variance. Short prompts, long generations, or `STRATA_QPN8=1` show the decode gains more clearly.

## Numerics summary

Bitwise: the operator fusions, the int8 KV gather, the NVLink P2P hand-off, QPN8's quantized intermediates. FP32-level (summation order differs; bounds in the headers and parity tests): Volta prompt attention, BF16→FP16 GEMM, GDN chain-split, grouped verify attention (≤1e-3× output scale). Every sm_75/80+ path is bit-for-bit unchanged.

## Tried and dropped (negative results, recorded so nobody repeats them)

- QPN8 enabled by default: windows below M=4 get slower (0.66–0.91×), so it stays opt-in.
- Chunked GDN (hand-written FlashQLA port): correct but 4–6× slower than the recurrence (128 threads/CTA cannot feed the tensor cores and every stage is latency-bound; llama.cpp-v100's +5.4% relies on TileLang-level pipelining). Kept opt-in.
- q4_0 KV: at 32k its gather is 2.3× slower than int8 (32.0 vs 13.9 µs) and it also quantizes K; k8v4 is worth it only when VRAM is the hard limit. V100/32k should stay on `--kv int8`.
- Tensor cores (m8n8k4) for decode GEMVs over DP4A: M=1 is bandwidth-bound and the tensor-core path is slower (same conclusion as bench/results/2026-09-30-tc-gemv-sm75 on Turing).
- `CUBLAS_COMPUTE_16F` for the BF16 GEMMs: rejected by cuBLAS on CUDA 12.8; `CUBLAS_GEMM_DEFAULT_TENSOR_OP` matches the default in speed, defaults unchanged.