#1123 HIP RDNA2 (gfx103x): prompt GEMMs as FP32 SGEMMs, a faster default that cannot overflow (RX 6800: 343 -> 534 tok/s)

Descripción

Resubmission of #1006, which GitHub closed when `main` was force-pushed. The branch is now on the new `main` (82f46a8), so it sits beside #835's `STRATA_HIP_PROMPT_F16`.

On gfx1030, the rocBLAS that setup installs (ROCm 10.2.0a20260930) has tuned kernels for FP16 -> FP16, int8 and FP32 -> FP32 only. The prompt path's dense GEMMs take FP16 or BF16 inputs and write FP32 (rocBLAS's HS / BS types), so they run fallback kernels at about 5 TFLOPS. That is why a default gfx103x install reads prompts at about half the speed `STRATA_HIP_PROMPT_F16=1` reaches.

This PR is meant as a faster default for gfx103x that does not depend on the FP16 range. It widens both inputs to FP32, which loses nothing, and runs SGEMM, which accumulates in FP32 like the native call. No value passes through FP16, so it cannot overflow. On the FP16 route, a GEMM output above 65504 becomes inf, and that can then spread through the layer as NaN. That is a hard failure rather than a rounding difference, and nothing bounds it for models that nobody has range-checked.

**What it is not:**
- It is slower than `STRATA_HIP_PROMPT_F16=1`.
- It is not bit-identical to the current default. The two differ by at most 4.3e-6 of the largest output on random inputs, about a hundredth of FP16's rounding step. Like any change of that kind, it can flip a near-tied greedy token.

If you would rather keep the default bit-identical, I can make the route opt-in instead. It would then mainly serve as the overflow-free alternative to the FP16 route.

| RX 6800, 12K-token prompt | tok/s |
|---|---:|
| `main` default (`STRATA_RDNA2_SGEMM=0`) | 343 |
| this PR (default on gfx103x) | 534 (+56%) |
| `STRATA_HIP_PROMPT_F16=1` | 619 (+81%) |

**Changes:**
- `Gemm::rdna2_sgemm` (`src/prefill/gemm.cu`) widens the weight once per call and the activations in slices of up to 128 MiB of FP32.
- **Order in `Gemm::bf16` / `Gemm::f16`:**
  1. the FP16 route, when `STRATA_HIP_PROMPT_F16=1`;
  2. hipBLASLt;
  3. this route;
  4. the native `hipblasGemmEx`.

  With the FP16 route on, this route only sees the products it leaves: `f16` with `beta != 0`.
- **When it runs:** it is on by default on gfx103x only. Other AMD architectures and CUDA are unchanged.
  - N < 64 (routers, gates) keeps the native call.
  - A padded BF16 X (`ldx`, `STRATA_PF_PAD`) keeps the native call.
  - If the FP32 buffers cannot be allocated, the native call runs.
- `STRATA_RDNA2_SGEMM=0` / `=1` turns it off / on for any AMD card.
- `docs/AMD_HIP.md`, RDNA2 section: the measurements, how the two routes combine, and a Windows gfx1030 report.

**Measured** on an RX 6800 16 GB over OCuLink (PCIe 4.0 x4, 7.1 GB/s), Ryzen 7 8845HS, 28.8 GB RAM, Windows 11.
- **Setup:** this branch built with `tools/hip/build_windows.bat` (`STRATA_HIP_ARCHS=gfx1030`); Coder IQ1_M with setup's arguments (64K context, MTP); fresh 12K-token prompts with no reuse.
- **Method:** each arm ran four prompts in one session with one binary. The table above averages the last three of each arm, because the first prompt after a cold start reads more from disk.
- **Per call** (the engine's exact `hipblasGemmEx`: opA = T, opB = N, FP32 out, fp32 compute):
  - N 10240 x T 8192 x K 2560: 86.7 -> 27.7 ms, including the widening.
  - N 320 x T 8192 x K 10240 (BF16): 10.9 -> 7.5 ms.
- **Numerics:** I compared the native call and this route with a float64 reference, using random inputs in [-1, 1] and 10 FP16 / BF16 shapes. This route was as close as the native call or closer on all 10. The point is only that nothing is lost; the difference from the native call is far below the BF16 rounding the prompt path's activations already carry.
- **FP16 range on this model:** `STRATA_F16_RANGE=1` on the Coder IQ1_M (four 12K prompts) peaked at 105.3 for activations and 424.8 for FP16 outputs, with nothing beyond 65504. So `STRATA_HIP_PROMPT_F16=1` is fine for this model. This route is for models nobody has checked, where an output past 65504 would become inf instead of a large finite number.
- **End to end:** a 9-task code-repair benchmark, where the model patches seeded bugs in a browser game until 38 browser checks pass. It solved 9/9 on the first try with native, SGEMM and FP16 alike. Model time was 458 s with native, 374 s with SGEMM and 307 s with FP16 (engine 0.1.39 + #1006 / #835).
- **Tests:** `tests/hip/prefill_gemm.cpp` gains two cases that cross the activation slice boundary (T 3300, K 10240): FP16 with `beta = 1` and a padded `ldy`, and BF16. `hip_prefill_gemm` and `hip_gemm_f16_io_parity` pass on the RX 6800.
- **Full ctest** (the RX 6800 with the 780M hidden by `HIP_VISIBLE_DEVICES=1`): 72 passed, 7 skipped and 7 failed out of 86. None of the failures goes through `Gemm`:
  - `ple_parity`, `expert_parity` and `pool_test` need model fixtures that are not on this PC.
  - `expert_cache_segmented_test`: `--vram-elastic` is CUDA-only.
  - `hip_handoff` and `hip_prefill_mmq_parity` also failed on the old `main`.
  - `hip_prefill_wmma_gemm_parity` (new on `main`) calls the WMMA kernel directly. Every case is declined on gfx1030, which has no WMMA, and the test counts that as a failure.
- **Another RX 6800:** in #1006, @did-technomancer measured the previous version against #835 on a second RX 6800 under Windows (IQ3_XXS, 128K context). Prompts were 14–16% slower with SGEMM than with FP16, and decode was the same. That matches the gap here (534 against 619, 14%). That comparison did not include the default path.
- **Not tested:** Linux, and other gfx103x cards. In #1006, xjc10 measured the previous version on an RX 6900 XT under Linux at +51 to +65%.

Thanks to @xjc10 for #835 and for the comparison on #1006.