#835 HIP gfx103x: the prompt path's 16-bit GEMMs in FP16 in and out (prompt reads ~2x on RX 6900 XT)

Description

## Summary

On gfx103x (RDNA2) rocBLAS has tuned GEMM kernels for FP16 in / **FP16 out** only. The prompt path's 16-bit products run
FP16 in / FP32 out (`Gemm::f16`, `Gemm::native`) and BF16 in / FP32 out (`Gemm::bf16`), which on gfx1030 fall back to generic
kernels: 5.6 and ~5.3 TFLOPS against 37.7 for FP16 out (N=10240, T=7313, K=2560). On a 7.3K prompt the FP16-in/FP32-out
fallback alone was 54% of the GPU time.

With this change, on gfx103x only:
- `Gemm::f16` / `Gemm::native` run FP16 out into the start of each FP32 row of Y (`ldc = 2·ldy` halves) and
  `widen_rows_f16` widens each row in place from its end: no extra buffer, no host sync.
- `Gemm::bf16` converts W in row slices to FP16 through the existing dequantization scratch (as `native()` does), and the
  kernels that write the BF16 GEMMs' activation images (`gr_*`, `to_bf16`) write FP16 instead (saturated like the
  existing FP16 images; a NaN stays a NaN). BF16X2 is off there.
- It is an explicit per-instance switch, `Gemm::set_f16_io`, turned on only by `Prefill::init` (with `set_act_f16`), so
  every other caller - e.g. `gemm_bf16_parity` - runs exactly what it ran. `prompt_f16()` is cached per device, since a
  layer split can mix cards. `STRATA_HIP_PROMPT_F16=0/1` overrides. CUDA builds are unchanged.

Not the same as #655 / #540 (FP16 tensor-core path below sm_80): that one keeps FP32 output, which on gfx1030 is the slow
case; HIP is untouched by it.

## Measurements (2x RX 6900 XT, ROCm 10.0, IQ3_S; details in `bench/results/2026-10-04-rdna2-fp16-prompt/README.md`)

| | old path | this change |
| --- | ---: | ---: |
| one card, 9.4K / 34.7K / 105.8K prompt | 439 / 466 / 461 tok/s | 744 / 915 / 926 tok/s |
| `--layer-split auto`, 8.3K / 33.6K | 444 / 707 tok/s | 833 / 1,356 tok/s |
| one card + expert helper, 8.3K / 33.6K | 432 / 454 tok/s | 798 / 887 tok/s |

Decode unchanged. Same binary, `STRATA_HIP_PROMPT_F16=0` vs default.

## Checks

- **Range:** over ~110K tokens of English docs, Chinese logs and C++, no BF16 GEMM input, weight or FP16 output beyond
  65504 and no infinity.
- **Needle** 8k/32k/128k × 5 depths: 15/15 on both paths.
- **Teacher-forced distribution** (the method of the V100 report, 3 chats, 1,578 scored positions): each path is
  bit-identical on a rerun. Against a higher-precision reference (old path with `STRATA_PREFILL_BF16X2=1`), the FP16 path
  is closer than the old path: median KL 0.0034 / 0.0044 / 0.0038 vs 0.0092 / 0.0096 / 0.0055, argmax agreement
  95.5 / 96.5 / 95.8% vs 93.1 / 93.3 / 93.8% (FP16 activations keep 10 mantissa bits where BF16 keeps 7).
- 17 server starts in these runs, no stall.

## Limits

One machine, one model (IQ3_S packed from a non-setup GGUF), one corpus. gfx1031/1032 share the `gfx103` prefix but were not run.

## Validation and docs

- **ctest** (`-DSTRATA_BUILD_TESTS=ON`, the same options as the engine build, gfx1030, ROCm 10.0): 60 of 65 pass, the same 60 as `6f32ec0` without this change. Skipped: `hip_prompt_attn_wmma` (no matrix cores on gfx1030), `hip_prefill_hipblaslt_gemm` (no hipBLASLt table). Failed for reasons outside the engine, identically with and without the change: `ple_parity` (a Q2_0 PLE file that is not on this machine), `expert_multi_test` (this Zen 3 CPU has no AVX-512), `platform_memory_test` (`mlock` at the shell's default `ulimit -l`).
- **Docs:** a bullet in `docs/AMD_HIP.md`'s RDNA2 section: what changes, the switch, the measured numbers and the ctest result above.

Developed with an AI coding assistant; every number above was measured on 2x RX 6900 XT (gfx1030; one card for the single-card rows) / Ryzen 5 5600X, ROCm 10.0.