Pull requests / #176

#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)

closed · @Efeisot · 0 comments · View on GitHub

BenchmarksSetup & installMulti-GPUAMD / HIPNVIDIA / CUDAModels & quantsDocumentationWindowsLinux

Description

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)

Related on strata.com

Editorial links to help you install, pick models, or read release notes — not part of the upstream thread.