#176 HIP backend on gfx1200 (RDNA4, RX 9060 XT): the gfx1100 backend runs unchanged with three deltas (Changes made by GLM5.3-Flash from freebuff)

Descripción

This carries the merged AMD HIP backend (PR #121, engines 0.1.25/0.1.26) to gfx1200 (RDNA4,
RX 9060 XT 16 GB, ROCm 7.2 / hipBLASLt 100202, wave32 verified at run time) - a card this
repository has not tested. The backend's assumptions hold on RDNA4: wave32 and the 64 KiB workgroup
LDS limit are unchanged, so the gfx1100 code runs with three deltas:
 
* cmake/hip_backend.cmake: the arch gate accepts gfx12xx beside gfx11xx.
* src/core/device.cu: the runtime device check accepts gfx1200 beside gfx1100.
* include/math_constants.h (new): a fallback for ROCm installs that do not ship
  <math_constants.h> (this 7.2 one). The hip_compat include dir is searched after the toolchain's,
  so a real header always wins; the engine references only CUDART_INF_F/NAN_F/PI_F, all defined.
 
One data file: tools/hip/gfx1200-hipblaslt-100202.txt, hipBLASLt solutions calibrated on the card
with tools/hip/tune_hipblaslt at the shapes the engine launches (the shipped gfx1100 tables do not
dispatch here: the runtime guards arch and version and falls back to plain hipBLAS). 7-15x over
plain hipBLAS per shape, +43% prefill end-to-end. The recalibration command for other cards is in
docs/AMD_HIP_GFX1200.md.
 
One bitwise-neutral tweak on the portable attention path (the sm80 tensor-core kernel is compiled
out under HIP, so prefill runs the decode kernel): #pragma unroll 4 on the value accumulation.
A second experiment - prefill query batch 32 -> 64 for L2 reuse of overlapping selections - measured
within noise on 0.1.26 and is NOT included: rebased onto 0.1.27 it overflows the prompt path's
borrowed buffer region ("device buffers for a chunk of 256 tokens do not fit" at any chunk). That
may be worth a look on your side: 0.1.27's lend math appears to have little headroom over the
buffers it counts (iq1_m, 16 GB card, --prefill 2560 and every smaller chunk).
 
Tested on the card, rebased onto 0.1.27 (a790805): complete HIP build, 30/30 ctest with the Lt table
(the two documented exclusions apply), deterministic greedy output across runs, 512-token and
back-to-back runs without stalls, 1.4 GiB VRAM free at steady state (--vram-reserve-mib 1792;
1024 left ~360 MiB with a desktop running). Measured, IQ1_M Coder, greedy, MTP:
 
  2,374-token prompt, 128 new:   31.0 tok/s decode   541 prefill
  65,045-token prompt, 512 new:  26.3 tok/s decode   761 prefill   (--kv int8 --kv-resident 65536 --prefill 16384)
  130,091-token prompt, 32 new:  22.7 tok/s decode   637 prefill
 
Rejected arms and numbers are recorded in bench/results/2026-09-30-gfx1200/README.md (k8v4 vs int8
streaming on 16 GB, pool-worker counts, chunk sizes, and gfx1200's multiProcessorCount reporting the
WGP count: doubling ggml's nsm was A/B-tested, 726.6 vs 731.1 tok/s at 64K, and not kept).
 
New files: docs/AMD_HIP_GFX1200.md (setup, tests, hipBLASLt recalibration, running, limits),
bench/results/2026-09-30-gfx1200/README.md, gfx1200-run.sh (launcher for the measured
configuration; model paths via env vars). gfx1200-run.sh is root-run-*.sh-ignored, hence the name.
 
Not validated here: other RDNA4 cards, Windows HIP, multi-GPU, vision, answer-quality benchmarks,
and full 262144-token contexts (tested to 131K). The gfx1100 caveats in AMD_HIP.md apply unchanged.
 
Note: I can't ssh my system to you even if i wanted due to some bad conditions, maybe the patches made by GLM5.3-Flash from freebuff can help you to find something that we didn't see, thanks for the effort
Tested System: CachyOS with Linux 7.2.8-1, RX9060XT Pure 16GB GPU with 64GB DDR5 RAM and R9 7950X CPU (tdp limited to 65w)