#2152 hip(RDNA4): a parity-gated fused-GR register-half down kernel for gfx1200/gfx1201 decode

Descrição

### Summary

Adds one extra selection for the hyper-connection fused-GR read on AMD RDNA4: the down-projection kernel
`gr_down_register_half_kernel` ("P4"), which carries the next activation tile in a compile-time register tuple
before scattering it into the existing shared-memory buffer. It is gated to runtime `gfx1200`/`gfx1201` only, and
it is selected **only when asked for** — `STRATA_HC_SPLIT=7` — and only after a per-device self-test proves it is
**bitwise identical** to the existing staged kernel. If any gate fails, the established staged / split / plain
path is selected unchanged.

This is not a replacement of the gfx12 kernel set and not a new arithmetic path. The accumulation order, tile
order, grid, block size and dynamic shared-memory layout are the staged kernel's; only the way the `h+2` tile
reaches shared memory differs.

Opt-in surface: `STRATA_HC_SPLIT` (new selector value `7`, existing `0`/`1`/`2`). **Unset stays on the staged
upstream path on every architecture: it never runs the P4 self-test and never selects P4.** `=2` requests staged,
`=1` split, `=0` plain. Only `=7` requests P4, and only on `gfx1200`/`gfx1201`, still subject to both parity
gates. Every other architecture keeps upstream's staged selection.

Files: `src/kernels/cuda/fused_gr.cu`, `include/strata/kernels/fused_gr.hpp`,
`include/strata/kernels/fused_gr_policy.hpp` (new), `tests/kernels/fused_gr_p4_policy_test.cpp` (new),
`CMakeLists.txt`, `docs/AMD_HIP.md`.

### Motivation

The fused-GR read is part of the per-token decode path and stages activation tiles for each hidden-state tile.
P4 tests whether an explicit per-thread register prefetch can hide some of that data-movement cost without
changing the dot products. The existing staged implementation fills a shared-memory double buffer through
`stage_htile`; P4 preserves that shared layout and the staged arithmetic order.

P4 splits the next tile into two compile-time halves: half A is already in registers before the tile is needed,
and half B is fetched in the middle of the dot-product loop. The compile-time `HcChain` avoids a runtime-indexed
per-thread array that the source comment says could lower to private scratch on gfx1201. Fetching half B at `q==2`
is intended to expose load/compute overlap with the remaining dot products; **that overlap has not been confirmed
by an ISA dump or profiler**. The arithmetic and q/k order are unchanged, so bitwise parity with staged is required.

Related prior work, distinct from this change: #2070 (CUDA sm_75 hyper-connection up kernel and token-pair
staged down kernel), #1370 (CUDA sm_80+ opt-in `STRATA_HC_UPMIX`), and closed #329 (gfx12 WMMA for QSA prompt
attention). None of them touch the RDNA4 fused-GR decode down path.

### Implementation

`src/kernels/cuda/fused_gr.cu`:

- `HcChain<I>` — a recursive struct of `float4` slots indexed only at compile time. The source comment states
  this is to keep the per-slot activations from lowering to private scratch on gfx1201; no ISA dump or profile
  is included with this PR.
- `gr_down_register_half_kernel<MAX_T, EXACT_T>` follows `gr_down_staged_kernel` exactly for tiles 0 and 1
  (`stage_htile` + `cp_async_commit`), then handles tile `h+2` in two halves per thread:
  - `NS = ceil(MAX_T * (H_TILE/4) / THREADS)` slots per thread, split at `H0 = (NS+1)/2`. For the generic
    instantiation (`MAX_T = 8`, `H_TILE = 1280`, `THREADS = 256`) that is 10 slots, 5 per half, i.e. 40 floats
    per thread in registers.
  - half A (`t0`) of tile `h+2` is loaded one iteration ahead — at the end of iteration `h` the kernel issues
    `hc_load_all` for tile `h+3`, which is iteration `h+1`'s half A (tile 2's half A is loaded before the loop);
  - half B (`t1`) of tile `h+2` is fetched inside the dot loop at `q == 2`, so that its global loads are issued
    while the remaining dot products of tile `h` are still executing. **This is the intended schedule, not a
    measured overlap** — no codegen or profiling evidence is attached.
  - both halves are then scattered into the shared buffer. `hc_scatter_all` uses the same index swizzle as
    `stage_htile` (`k*(H_TILE/4) + (s4&1)*(H_TILE/8) + (s4>>1)`), so the dot products read the same shared
    layout, and the same bits.
  - Inactive lanes (the inject block's non-participating warps) skip the mid-dot prefetch and perform their own
    single fetch of half B, because they still take part in the shared-buffer scatter.
- Launch shares staged's grid (`DOWN_BLOCKS + 1`, 256 threads) and dynamic shared-memory size, and the kernel
  instantiations get the same `cudaFuncSetAttribute` treatment as the staged ones (max dynamic shared memory,
  100% shared-memory carveout).
- `fused_gr_read_multi` dispatches `kHcRegisterHalf` (=7) to this kernel; plain/split/staged dispatch is
  unchanged.
- `fused_gr_check()` and `fused_gr_variant()` ask the policy header what the environment asks for, so the
  selector rules live in one place and are unit-tested. `fused_gr_variant()` reports the requested path capped at
  staged until the check has run, so an untested `=7` request never reaches the candidate kernel.

`include/strata/kernels/fused_gr_policy.hpp` (new, header-only, no GPU) holds the selection and error rules as
`constexpr` functions so they are unit-testable:

```cpp
constexpr int fused_gr_env_variant(bool env_present, char first) noexcept;     // unset -> staged
constexpr bool fused_gr_p4_selftest_requested(int want, bool gfx12) noexcept;  // only =7, only on gfx12
constexpr int fused_gr_select(int want, bool split_ok, bool staged_ok, bool p4_gate_ok) noexcept;
constexpr bool fused_gr_p4_parity_gate(bool gfx12, bool staged_ok, bool p4_ok) noexcept;
constexpr bool fused_gr_p4_q32_parity_marker_ok(bool selftest_ok, bool qfuse_hc, bool staged_ok, bool p4_ok) noexcept;
enum class GrBackend : unsigned char { Cuda, Hip };
constexpr bool fused_gr_context_fatal(int err, GrBackend backend) noexcept;
```

`fused_gr.cu` `static_assert`s every code the classifier uses against its own backend's header name
(`cudaError*` in `driver_types.h`, `hipError*` in `hip_runtime_api.h`) in the backend being compiled, so a header
that renames or renumbers one breaks the build instead of the classification. The CUDA-only codes (226, 714–718,
721, 810, 911, 912) are asserted in the CUDA branch and the HIP-only HSA runtime codes (1052, 1053) in the HIP
branch; HIP's different meanings for 911 (`hipErrorInvalidChannelDescriptor`) and 912 (`hipErrorInvalidTexture`)
are asserted where they differ, together with the classification those numbers get in that backend.

### Correctness and fallback

Architecture gate. `p4_gfx12_device()` reads `cudaDeviceProp::gcnArchName` at runtime and accepts only
`gfx1200` or `gfx1201` (`gfx_arch_is_gfx12`). Under a non-HIP compile it returns false. An explicit
`STRATA_HC_SPLIT=7` on any other part logs that P4 is restricted to gfx1200/gfx1201 and keeps staged.

Opt-in gate. The P4 self-test runs only for an explicit `=7` on gfx1200/gfx1201. Unset, `=2`, `=1` and `=0` never
run it, and the selector never promotes P4 on their behalf: unset with every gate passing stays staged. On
gfx1200/gfx1201 with anything but `=7` the check logs that P4 is opt-in through `=7`.

Promotion gate. P4 is selected only when `gfx12 && staged_ok && p4_ok` — that is, the staged kernel must itself
match the plain reference bitwise, **and** P4 must match staged bitwise. A P4 candidate that passes while the
staged reference is failing cannot be promoted.

Self-test coverage (`fused_gr_selftest`, run once per card and latched in `fused_gr_check`): T = 1..8, apply
off/on, injection weights present/null, and with and without the q8_1 path (the q8 arm also exercises the
`QFUSE_HC` q32 up kernel). Compared bitwise: `lo`, `rs`, `inject`, `mixed`, `R_out` floats **and** the `q8_1`
output bytes.

Failure isolation, split by what the error leaves usable. The candidate runs on its own non-blocking stream,
which isolates **recoverable candidate-only** errors: a refused P4 launch, a bad configuration/handle, P4 stream
creation, final synchronization, or stream-destruction failure disables P4 without invalidating staged/split checks.
Shared scratch-allocation, reference-stream setup/sync/destruction, or shared cleanup failures fail the self-test, so
P4 cannot be selected. A separate stream does **not** isolate an asynchronous device fault — an illegal address or
device exception from the candidate poisons the shared context for the whole process, so fallback kernels would only
return the same error. Context-fatal errors returned by device/property/attribute queries, H2D copies or memsets,
shared-memory setup, launch/sync, read-back, allocation, stream creation/destruction, final sync or cleanup are
classified before clearing an error or making another GPU call; they log the code, flush output and call
`std::_Exit(1)` immediately. This matches #224's handling in `src/program/generate.cpp` and avoids unwinding/
destructors on a poisoned context. The fatal set includes illegal address (700), timeout (702), destroyed context
(709), device assert (710), launch failure (719), and uncorrectable ECC (214) in both backends, plus CUDA's 226,
714–718, 721, 810 and 911. HIP's 911/912 are known recoverable API errors; HIP's 1052/1053 are documented HSA-call
failures. Unknown (`cudaErrorUnknown`/`hipErrorUnknown`) and all other unclassified runtime values fail closed; only
explicitly listed ordinary API/resource errors (invalid value/configuration, allocation/launch-resource refusal,
invalid device function/image, bad handle, not-ready, unsupported limit/capture, etc.) are allowed to fall back.

Selector semantics (`STRATA_HC_SPLIT`): **unset = staged, the upstream path** — no P4 self-test, no P4 selection;
`=2` requests staged; `=1` selects split; `=0` selects plain; `=7` is the only value that requests the P4
self-test, on gfx1200/gfx1201 only, and selects P4 only if both gates pass, falling back staged → split → plain
otherwise.

Diagnostic marker. The `P4 + QFUSE_HC q32 outputs passed bitwise parity` line is emitted only when the whole
self-test succeeded, `QFUSE_HC` is enabled, and both parity flags passed — it cannot print after an unrelated
self-test failure.

### Measurements

These are previous same-host results, not a new run on the rebased branch. The host was a Radeon AI PRO R9700
(native gfx1201), ROCm 10.2.0 nightly / HIP 7.17.26391 / hipBLASLt 100500, running Swift / Qwen3.8-Flash-Next
IQ3_XXS. Each arm used a 439-token warmup and three fresh 80-line / 1,111-token outputs at `temperature=0` and
`cache_n=0`. The staged and P4 arms used the same v0.1.43 serve module and **same integration binary**, serially,
differing only by selector. That binary also carried independent local MMQ/hipBLASLt overlays, held constant across
both arms; these numbers are an integration-build A/B, not a clean-upstream benchmark.

| Arm | Median decode | Range (3 trials) |
|---|---:|---:|
| v0.1.43 staged upstream path (`STRATA_HC_SPLIT=2`) | 153.2 tok/s | 152.2 – 154.3 |
| v0.1.43 P4, that build's promoted default (unset) | 156.1 tok/s | 153.3 – 156.8 |

P4's observed median is **+1.89% vs the staged upstream path**. The ranges overlap (153.3–154.3), so this is not a
statistically established gain: three trials on one host/workload. **The 156.1 tok/s P4 timing comes from the
earlier integration build, where an unset selector promoted P4 once both gates passed.** The explicit `=7`
selector that this update makes the only P4 request was exercised for correctness in those runs, not for timing,
and no P4 or clean-upstream decode run was performed on the rebased v0.1.43.1 tree because the live service was
kept online.

Correctness observed in the existing runs: P4 and staged replies were byte-identical in all three paired trials.
The four tested selector configurations — `=2` staged, `=1` split, `=7` P4 opt-in, and that build's unset P4
default — all returned the same arithmetic answer `391`