Pull requests / #835
#835 HIP gfx103x: the prompt path's 16-bit GEMMs in FP16 in and out (prompt reads ~2x on RX 6900 XT)
closed · @xjc10 · 0 Kommentare · Auf GitHub
BenchmarksSetup & installMulti-GPUAMD / HIPNVIDIA / CUDAModels & quantsDocumentation
Beschreibung
## 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.
Mehr auf der Site
Links zu Install, Modellen, Releases.