#315 hc: the multi-token hyper-connection read split finer

描述

The multi-token hyper-connection read (`fused_gr_read_multi`: every verify window and the MTP drafter) keeps v1's arithmetic exactly - every sum over the same elements in the same order - and only splits the work differently. On an RTX 4080 SUPER it takes the two HC reads of a verify window from 5.01 to 4.21 ms (-16 %), about 0.9 ms of GPU time per round, with the same tokens.

## What changes

- **Norm (v2, v3):** one block per token *and stream* instead of one per token. Each thread visits the elements of its stream in the order it did in v1, so every sum of squares and `rs` is v1's; `R' * w_norm` is recomputed and scaled once (the same two roundings as v1's store and in-place scale).
- **Down projection (v3):** v1's rows per block and lane order, but the activations arrive by `cp.async` in half-stream tiles (1280 floats: 160 = 5 x 32 chunks, so a lane still accumulates its chunks `lane + 32 q` in ascending order), two in flight, in a bank-conflict-free two-plane layout, and the next tile's weights load while the current one is used. Before sm_80 and on HIP the staging is a plain copy (the same bits).
- The down kernels carry as many tokens per launch as the card's opt-in shared memory allows (v3: 10 KB per token), sliced like v1 since 0.1.27; pre-Volta cards use the per-block limit.

## Exactness

- `Verifier::init` runs `fused_gr_check()` once per card: v1, v2, v3 and the single-token read on random bf16 weights and inputs (1..8 tokens, with and without the pending write), every output compared bit for bit. The newest variant that agrees runs on that card, and the log says which:
  `strata hc: CUDA0: the hyper-connection read runs as v3 (...); checked bit for bit against v1 on this card`.
  A card that has not been checked runs v1. `STRATA_HC_V2=0` keeps v1, `=2` stops at v2.
- End to end: requests without anything timing-dependent (no prompt-lookup drafts, no expert swaps, `pcie_frac=0`) give the same 256 tokens with v3 and v1, greedy and sampled (T 1.0, top_p 0.95, top_k 20, seed 1), in the same number of rounds.

It is on by default where the check passes, since it changes no bit. If you prefer it opt-in, that is one line in `env_variant()`.

## Measured

RTX 4080 SUPER 32 GB (sm_89), Ryzen 7 5800X3D, 64 GB DDR4, PCIe 3.0 x16, Windows 11, CUDA 13.3; IQ3_S, `--max-context 262144`, int8 KV, 11055 experts in VRAM. A local test build of 0.1.30 (30ec18e) with this commit, #242, #284, and #186 and #258 behind switches; every comparison is within that one build (`STRATA_HC_V2` unset = v3 against `STRATA_HC_V2=0` = v1), so everything else is identical.

The deterministic requests above (8K context, 256 tokens, the same tokens and 100 / 99 rounds in both):

| | v1 | v3 | |
|---|---|---|---|
| greedy: ms per round | 26.43 | 25.71 | -2.7 % |
| sampled: ms per round | 25.00 | 24.28 | -2.9 % |
| GPU wait per round (greedy / sampled) | 12.08 / 12.46 | 11.20 / 11.57 | -0.88 ms |
| tokens/s (greedy / sampled) | 96.9 / 103.4 | 99.6 / 106.5 | +2.8 % / +3.0 % |

With prompt-lookup drafts and expert swaps (3 x 256 tokens per cell, 8K and 64K context, greedy and sampled) tokens/s scatter by several percent from request to request, but the GPU wait per round drops in every cell: 14.67 -> 13.63, 14.90 -> 13.77, 15.94 -> 14.91, 16.00 -> 15.03 ms.

GPU stages (`STRATA_VERIFY_PROFILE=1`, 8K, ms per verify window, DeltaNet + attention layers summed; the stamps split only the first read of each layer):

| | v1 | v2 | v3 | #186 (`STRATA_GR_V3=1`) | v1 + #258's 1280 tile, forced on sm_89 |
|---|---|---|---|---|---|
| read 1: norm | 0.55 | 0.42 | 0.42 | 0.27 | 0.55 |
| read 1: down | 1.02 | 1.02 | 0.76 | 1.16 | 0.92 |
| read 1: up | 0.64 | 0.64 | 0.63 | 0.68 | 0.64 |
| read 2 + router | 2.80 | 2.65 | 2.40 | 2.70 | 2.70 |
| **sum** | **5.01** | **4.73** | **4.21** | **4.81** | **4.81** |

On this card v3 is 0.6 ms per window faster than #186, and it keeps v1's bits.

To reproduce on the same build: `STRATA_HC_V2=0` (v1) against unset (v3); `STRATA_VERIFY_PROFILE=1` for the stages.

Tested on sm_89 only. The pre-sm_80 and HIP paths compile the same kernel with plain copies instead of `cp.async`; there the startup check decides as well.