| # AMD Radeon: the HIP backend (gfx1100, gfx1101, gfx1200, gfx1201, gfx1030) |
|
|
| Strata runs on AMD Radeon cards through its HIP backend, the same engine as on NVIDIA compiled for AMD. This page |
| covers the build on Linux (on Windows a ready-made engine, see [Windows](#windows)) for the RX 7900 XT / XTX (RDNA3, gfx1100) and the |
| RX 9070 / 9070 XT / Radeon AI PRO R9700 (RDNA4, gfx1201; see [RDNA4](#rdna4-gfx1201)). The RX 7800 XT / 7700 XT |
| (gfx1101) and the RX 9060 XT (gfx1200) were validated by their owners (see [Community-validated |
| cards](#community-validated-cards)); the RX 6800 / 6900 series (RDNA2, gfx1030) builds and runs too, reported by a community machine and not yet validated by the maintainers (see [RDNA2](#rdna2-gfx1030)). Setup chooses it by itself on a PC with no NVIDIA card Strata can use (`--backend hip` on a PC with both); the |
| install steps for users are in [INSTALL.md](INSTALL.md#amd-cards). Other AMD architectures, wave64, and mixed |
| AMD/NVIDIA execution in one run are not supported. |
|
|
| The backend maps the CUDA-shaped runtime and BLAS calls to HIP/hipBLAS, uses |
| RDNA2/RDNA3/RDNA4's signed integer dot instruction for quantized kernels, and supplies |
| wave32 shuffle/packed-byte operations. CUDA-only QSA matrix instructions have |
| an ordered FP32 fallback. Prefill supports both dequantization plus hipBLAS GEMM and opt-in HIP ggml MMQ. |
| An optional, calibrated hipBLASLt path accelerates dense projections (per-architecture tables in `tools/hip`). |
| This does not claim bit-identical model answers across backends. See |
| [performance settings and evidence](AMD_HIP_PERFORMANCE.md). |
|
|
| ## Install with setup (recommended) |
|
|
| On Linux with an RX 7900 XT / XTX, RX 7800 XT / 7700 XT, RX 9060 XT, RX 9070 / 9070 XT or Radeon AI PRO R9700 and |
| the kernel's amdgpu driver (no ROCm install needed): |
|
|
| ```sh |
| ./setup.sh --backend hip |
| ``` |
|
|
| - **Detection:** setup finds the card through the kernel's KFD topology. Integrated Radeon GPUs are listed as not |
| supported. On a PC without an NVIDIA card Strata can use, `--backend hip` is chosen automatically. |
| - **ROCm:** a system ROCm 7 in `/opt/rocm` (or `$ROCM_PATH`) with hipcc and hipBLAS is used when present. Otherwise |
| (or when it is older than 7.0) ROCm is installed into `.venv` from AMD's TheRock wheels (~10 GB, no sudo), pinned |
| to the version this backend was tested with, from the card family's index: `gfx110X-dgpu` for gfx1100 / gfx1101, |
| `gfx120X-all` for gfx1200 / gfx1201, `gfx103X-all` for gfx1030 (`STRATA_ROCM_VERSION` / |
| `STRATA_ROCM_INDEX` override them; the gfx1030 index is not checked to carry the pinned version: a system |
| ROCm 7 is the tested path there). |
| - **Engine:** compiled on your PC for the card's architecture (10-20 minutes, once; again after a `git pull` that |
| changes it, or when you pick a card of another architecture). This needs a C++ compiler and git |
| (`sudo apt install build-essential git`). |
| - **hipBLASLt tuning table:** setup uses `tools/hip/<arch>-hipblaslt-<version>.txt` only when it matches both the |
| card's architecture and the installed hipBLASLt version (read from `hipblaslt-version.h`; 1.2.0 is `100200`). |
| Otherwise it says so and the prompt's dense matrix products use plain hipBLAS (slower prompts, same answers). |
| A table's solution ids are valid only for that pair, and the engine refuses any other table. |
| - **Several cards:** setup takes one card (the one with the most VRAM, or `--gpu N`) unless you name more: |
| `./setup.sh --backend hip --gpus 1,0` splits the model's layers across them, the first one the main card (numbers |
| as setup lists them; `--gpus all` = every supported card, the most VRAM first). Every chosen card must be one of the |
| architectures above; the engine is compiled for each of them (cards of two families, e.g. gfx1100 + gfx1201, need |
| a system ROCm 7: AMD's wheels hold one family). A split pays only when no single card holds the model's experts |
| (see RDNA4 below). |
| - **Limits for now:** images only through the CPU encoder (`--vision cpu`, 0.1.32), no calibration. The Monitor |
| shows the card's load, VRAM, temperature and power from Linux sysfs (0.1.32). |
|
|
| The rest of setup is the same as on NVIDIA: the model download, the start script, the server. |
|
|
| ## Windows |
|
|
| Since 0.1.34 an AMD card on Windows is set up like an NVIDIA one: download Strata, double-click `START-HERE.bat`. |
| On a PC with no NVIDIA card Strata can use, the AMD card is chosen by itself; with both, setup asks |
| (`START-HERE.bat --backend hip` picks AMD directly). |
|
|
| - **You need:** Windows 10 or 11 (64-bit), one of the cards above, and a current AMD driver ([AMD Software: |
| Adrenalin Edition](https://www.amd.com/en/support/download/drivers.html)). Nothing else: no ROCm or HIP SDK |
| install, no compiler, no admin rights. |
| - **Detection:** setup reads the display adapters Windows lists (their PCI ids; the VRAM size from the display |
| driver's registry entry). An integrated Radeon is listed as not supported. |
| - **Engine:** the ready-made `strata-windows-x64-hip.zip` from the release (built by `tools\hip\build_windows.bat` |
| for gfx1100, gfx1101, gfx1102, gfx1200, gfx1201 and gfx1030) goes into `engine\`. It carries the ROCm libraries the |
| engine loads (`engine\rocm\bin`: the HIP runtime, hipBLAS / rocBLAS / hipBLASLt with their kernels for these cards, |
| amd_comgr and the Microsoft C++ runtime; ROCm 10.2.0a20260930 from AMD's TheRock builds, licenses in |
| `engine\rocm\licenses`). The HIP runtime works through the AMD driver's own components, so the driver is the one |
| thing it needs from the PC. |
| - **The HIP runtime next to `strata.exe` (0.1.35, #468 #461):** `amdhip64_7.dll` and `amd_comgr.dll` are also put in |
| `engine\` (setup copies them there on every start). Windows looks in the program's folder before System32, where |
| some AMD drivers install their own `amdhip64_7.dll`; with that one, the bundled libraries crashed on the first |
| prompt (an access violation, or `hipErrorInvalidDeviceFunction`). The engine's log names the runtime it loaded |
| (`strata generate: HIP runtime ...`). |
| - **Before the ~60 GB model download** setup runs `engine\strata-device.exe --list-devices` (with `engine\rocm\bin` on |
| the PATH): if the HIP runtime does not see the card, setup stops there and points to the driver. It also gives the |
| card's HIP number: with an integrated Radeon that is device 1, not 0 (#325). From then on setup lists the AMD cards |
| as HIP numbers them, so `--gpu N` and the config's `"gpu"` are HIP numbers. |
| - **Differences from Windows-on-NVIDIA and Linux-on-AMD:** no images yet (the CPU image encoder is Linux-only for |
| now), one card per model (`--gpus` is Linux-only for now), no calibration. |
| - Two Windows-only engine details (#247, #325): hipBLAS can return success and still leave `hipErrorInvalidValue` |
| set after some BF16/FP16 GEMMs (seen on gfx1201); the engine clears that one stale error after a GEMM that |
| succeeded, on Windows only. `hipHostGetDevicePointer` returns the host pointer itself on Windows: kernels read |
| mapped memory through it correctly, but a device-to-device copy into it does not land, which is why |
| `tests/hip/handoff` times out there (the engine does not use that copy; `tests/hip/mapped_alias` reports it). |
| - Two more (#380, #377, by BlueKingMuch, measured on an RX 6800 that drives the desktop): `hipMemGetInfo` on Windows |
| does not subtract what the desktop and other programs hold on the card, so `--expert-cache auto` filled the card |
| past what Windows keeps in VRAM and decode fell from 41 to 30 tok/s. The engine now lowers that free figure by |
| what Windows' video memory budget for the process withholds (logged once: `strata: Windows budgets N of this |
| card's M MiB ...`; `STRATA_WDDM_BUDGET=0` turns it off). And the PCIe probe times its copies on the host clock |
| there, since HIP's events read impossible speeds (3,300-26,000 GB/s), so a slow link now gets a smaller |
| `pcie_frac` as on NVIDIA. |
|
|
| **What is validated (0.1.34):** #325's author ran the engine of this port on an RX 9070 XT (Windows 11, ROCm |
| 10.2.0a20260930 in `.venv`, compiled on the PC): Coder IQ1_M at 32K, 29.2 tok/s decode, ~181 tok/s prefill, |
| correct answers; ctest 42 of 46. The maintainers have no Windows AMD card: the release zip was built on an NVIDIA PC, |
| and checked there on the Ryzen CPU's integrated Radeon (gfx1036, a test build of the same tree): with only the zip's |
| libraries on the PATH, `strata-device` lists the card and a hipBLAS BF16 GEMM matches the CPU; the HIP ctest passes |
| 52 of 56, the 4 failures the same as on the RX 9070 XT (`hip_handoff`, and three tests that need a pack fixture). |
| The ready-made zip itself has not run a model on a discrete card yet - please report. |
|
|
| **Reporting a Windows AMD run** (an issue, or on #325): your card and driver version (AMD Software > System), then |
|
|
| ```bat |
| engine\strata-device.exe --list-devices |
| engine\strata-device.exe --selftest |
| ``` |
|
|
| (from the Strata folder, after `set PATH=%CD%\engine\rocm\bin;%PATH%`), the end of `strata-<model>.log`, and the |
| speed lines the server window prints for a first answer. |
|
|
| **Building it yourself:** `tools\hip\build_windows.bat` (Visual Studio 2022 Build Tools with the C++ workload, Python, |
| git; no admin, no AMD GPU) installs ROCm from AMD's TheRock wheels into `.rocm-win`, builds, and packages |
| `dist\strata-windows-x64-hip.zip`; `START-HERE.bat --backend hip --prebuilt dist\` installs that one. |
| `tools\hip\build_windows.bat tests` also builds the HIP tests (`ctest` in `build-hip-win`, with |
| `.rocm-win\Lib\site-packages\_rocm_sdk_devel\bin` on the PATH). `STRATA_HIP_ARCHS` picks other architectures. |
|
|
| ## Build |
|
|
| Requirements: a working ROCm driver/runtime, HIP development headers and |
| compiler, hipBLAS and (for tuned dense prefill) hipBLASLt development files, CMake 3.24+, a C++20 host compiler, and Git. |
| The model still needs sufficient system RAM and fast SSD storage for its PLE |
| table. VRAM occupancy alone is not a throughput measurement. |
|
|
| ```sh |
| cmake -S . -B build-hip \ |
| -DCMAKE_BUILD_TYPE=Release \ |
| -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF \ |
| -DCMAKE_HIP_ARCHITECTURES=gfx1100 |
| cmake --build build-hip --target strata -j2 |
| ``` |
|
|
| `CMAKE_HIP_ARCHITECTURES` is `gfx1100`, `gfx1101`, `gfx1200`, `gfx1201`, or a list such as `"gfx1100;gfx1201"` |
| (one binary for both). gfx1102 (the same wave32, 64 KiB LDS and dot4 instruction) builds with a warning: it passed |
| ctest (#192) but no model run has been reported; so does gfx1030 (RDNA2: the older `v_dot4_i32_i8`, a community run in #311). At startup the engine and `strata-device` compare each GPU they use |
| (`gcnArchName` up to the `:` feature suffix) with the architectures the binary was compiled for, and require |
| wave32. A binary carried to another card stops with the card's name, its architecture and the build's list, |
| instead of failing later with "invalid device function". |
|
|
| If CMake cannot find the HIP compiler, add |
| `-DCMAKE_HIP_COMPILER=/path/to/rocm/llvm/bin/clang++`. On the Fedora-family test |
| host this was `/usr/lib64/rocm/llvm/bin/clang++`. Use the compiler and libraries |
| from the same ROCm installation. Do not enable both GPU backends in one build. |
|
|
| Native IQ experts are enabled by default. CMake fetches the llama.cpp revision |
| pinned by this repository. For an offline build, point `STRATA_GGML_DIR` at a |
| checkout of that exact revision; using an arbitrary newer checkout changes the |
| dependency being tested. |
|
|
| ## Model and serving configuration |
|
|
| Prepare a supported model pack and its MTP runtime using the existing tools. |
| For OrcaRouter IQ3_XXS, follow [the explicit compatibility conversion](ORCA.md); |
| this backend does not change quantization, model licenses, or tokenizers. |
| Add `--experts-bin` to the `tools/iq_pack.py` command: mmap requires the pack's |
| `experts.bin`, which the default native packing command does not emit. This |
| consumes additional disk space. Original GSQ-RCO IQ3_XXS uses shard 2 for PLE; |
| Orca uses shard 1. Keep each model's own pack and tokenizer together. |
| |
| Use `build-hip/strata` as the executable in the server JSON. Select the AMD |
| device with `ROCR_VISIBLE_DEVICES`/`HIP_VISIBLE_DEVICES` if necessary. Start with |
| a modest context and prefill chunk size before measuring larger workloads. |
| Keep the PLE file on an SSD. Do not assume a configured context length proves |
| successful full-window inference. |
| |
| Large pinned host allocations can fail on ROCm even when ordinary RAM is |
| available. The existing `--mmap-experts` path avoids allocating the full pinned |
| expert arena; it still depends on OS file-cache residency and may stall on |
| storage reads. It does not make SSD access equivalent to RAM. |
| |
| Starting args for the original GSQ-RCO IQ3_XXS model (replace the paths): |
|
|
| ```sh |
| build-hip/strata --serve \ |
| --pack packs/iq3xxs \ |
| --native /path/to/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00001-of-00002.gguf \ |
| --ple-gguf /path/to/Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS-00002-of-00002.gguf \ |
| --mmap-experts --expert-profile data/expert-profile.bin --expert-cache auto \ |
| --prefill 512 --spec 4 --spec-min-p 0.5 --mtp mtp/rt \ |
| --max-context 4096 --kv int8 --pool-workers 15 \ |
| --adapt-every 0 --pcie-frac 0 --vram-reserve-mib 1024 |
| ``` |
|
|
| `--serve` is the engine's internal token protocol. For a browser or OpenAI API, |
| put these args in the `args` array of the [server JSON example](ORCA.md#local-server), |
| set `exe` to `build-hip/strata`, and use the matching pack's `tokenizer` directory. |
| Launch with `.venv/bin/python -m serve.server --engine strata --config strata-hip.json --port 8080`. |
| The worker count above was used on a 16-core CPU; measure it for your CPU. |
| The 4K context is a smoke-test starting point, not a model limit. The expert cache |
| sizes itself automatically and leaves 1 GiB of VRAM headroom. |
|
|
| **The card also drives a Linux desktop (#560 #516):** keep more VRAM free than the default 700 MiB, e.g. |
| `./setup.sh --vram-reserve-mib 3072`. With the cache filling the card, the desktop's next VRAM need makes amdgpu move |
| GPU memory to system RAM (GTT), the OOM killer then ends KWin/plasmashell or `systemd-oomd` ends apps, or the |
| compositor fails with "Failed to pin framebuffer with error -12". |
|
|
| The installer supports this backend (see "Install with setup" above). Images run through the CPU encoder for now (`--vision cpu`). |
| Setup installs one AMD card, or several with `--gpus` (the engine's layer split; see RDNA4 below). |
|
|
| ## RDNA4 (gfx1201) |
|
|
| The RX 9070 / 9070 XT and the Radeon AI PRO R9700 run the same kernels as gfx1100: wave32, 64 KiB of LDS per |
| workgroup, the signed dot4 instruction (`v_dot4_i32_iu8` through `__builtin_amdgcn_sudot4`) and a 100 MHz wall |
| clock. The first report and patch came from doplxyz (#178). Validated on 2026-09-30 on doplxyz's test machine: |
| an RX 9070 XT 16 GB and a Radeon AI PRO R9700 32 GB (both gfx1201), a Ryzen 9 3900X (16 threads, no AVX-512), |
| 47 GB RAM, Ubuntu 24.04 in a KVM/VFIO guest. |
|
|
| - **Build:** complete HIP build with tests (`-DCMAKE_HIP_ARCHITECTURES=gfx1201 -DSTRATA_BUILD_TESTS=ON |
| -DSTRATA_PREFILL_MMQ=ON`), with the system ROCm 7.14 (hipBLASLt 1.4.1) and with setup's own path: TheRock |
| 7.10.0a20251120 wheels from `gfx120X-all` (hipBLASLt 1.2.0) and `build_engine_hip`. |
| - **ctest** (all 45 registered tests, R9700): 42 pass, `hip_prefill_hipblaslt_gemm` skips (no table), and two |
| fail for reasons outside the GPU: `ple_parity` needs the Q2_0 model fixture, `expert_multi_test` needs an |
| AVX-512 CPU. `hip_handoff` needed the volatile ring store: on gfx1201 a plain store to mapped pinned memory |
| stays in the GPU's L2 until the stream is synchronized. |
| - **Arch check:** a gfx1100-only build stops on the gfx1201 card at startup with the message above (engine and |
| `strata-device`). |
| - **End to end** (Coder IQ1_M, setup's arguments: `--expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5`, |
| MTP, `--kv int8`, `--max-context 32768`; 128 greedy tokens, system ROCm 7.14, no hipBLASLt table). The answers |
| are coherent and the same on both cards. The live server smoke (`serve/server.py`: web page, models, props, |
| chat, prefix reuse, streaming, the Anthropic endpoint, a sampled code answer) passed 9/9 on each card. |
| |
| | card | expert cache | peak VRAM | 4K prompt | 4K decode | 16K prompt | 16K decode | |
| |---|---|---|---|---|---|---| |
| | Radeon AI PRO R9700 32 GB | 12,288 slots, 23.4 GiB | 29.7 GiB | 982 tok/s | 45.5 tok/s | 1,402 tok/s | 48.3 tok/s | |
| | RX 9070 XT 16 GB | 4,931 slots, 9.4 GiB | 15.6 GiB | 782 tok/s | 30.8 tok/s | 1,235 tok/s | 35.4 tok/s | |
| | R9700, TheRock 7.10 wheels | 12,288 slots | | 957 tok/s | 44.2 tok/s | 1,284 tok/s | 43.9 tok/s | |
| |
| The engine's resident memory was about 26 GB in every run. Since 0.1.31 `__byte_perm` is one `v_perm_b32` and the |
| packed byte subtracts/compare work on four lanes at once (#262, ttio2tech): decode +15% on the R9700 (46.0 -> 53.0 |
| tok/s on a 4K prompt, 52.0 -> 60.5 warm) and +5-7% on the 9070 XT, prompts unchanged, the same tokens. |
| - **hipBLASLt:** the validation above used hipBLASLt 1.4.1. A table calibrated on the R9700 at the engine's shapes |
| with that version (0.98-1.76x per GEMM over hipBLAS) changed the end-to-end prompt speed by 0-3%, within noise, |
| so none was shipped for it: there the plain hipBLAS path is already close. For hipBLASLt 1.5.0 (ROCm |
| 10.2.0a nightly) `tools/hip/gfx1201-hipblaslt-100500.txt` is shipped (see "Tuning table" below); on one R9700 it |
| measured +3.9% prompt speed on 4,210-token prompts (1,590 vs 1,531 tok/s), a modest gain. |
| With the hipBLASLt 1.2.2 of a system ROCm 7.2.4 the plain path is far off, and `tools/hip/gfx1201-hipblaslt-100202.txt` |
| is shipped for it (24 of the gfx1100 table's 26 shapes; setup uses it only with that exact version): on an R9700 |
| with the full IQ3_XXS, `--kv int8 --kv-resident 32768`, a fresh 32K prompt read at 638 -> 1,177 tok/s and a 7K one |
| at 656 -> 1,164 tok/s with it, decode unchanged. `hip_prefill_hipblaslt_gemm` passes with it. |
| - **Both cards in one run (layer split, engine 0.1.30):** the config's `"backend": "hip", "gpu": [1, 0]` (R9700 |
| first) runs through `serve/server.py` (setup writes it with `--gpus 1,0` since 0.1.31). Auto split put layers 0-27 on the R9700 and |
| 28-47 on the 9070 XT. With every expert on the GPUs the split gives exactly the tokens of the R9700 alone (4K and |
| 16K prompts); checkpoints on the split (second turn, rewind, a prompt sharing a prefix, a cancelled prompt |
| retried) give exactly the tokens of a fresh read. The conversation cache refuses a split at start (exit 2). |
| On this pair the split does not pay: the R9700 alone already holds all 12,288 expert pairs. |
| |
| | run (the same session, warm) | 4K prompt | 16K prompt | decode | |
| |---|---|---|---| |
| | R9700 alone | 1,794 tok/s | 1,804 tok/s | 51-52 tok/s | |
| | R9700 + 9070 XT, auto split | 1,384 tok/s | 1,906 tok/s | 42-43 tok/s | |
| | RX 9070 XT alone | 1,016 tok/s | 1,519 tok/s | 38-40 tok/s | |
| |
| A split is worth it when no single card holds the model's experts. Since 0.1.31 a split pins the whole expert |
| arena on Linux (the 8 GiB cap is for Windows/WSL2 only, #253): with an expert cache of 1,500 on the R9700 (most |
| experts streamed) the split read a 16K prompt at 1,803 tok/s instead of 1,287, with the same tokens in 5 + 5 |
| starts. `"split_skip_if_fits": true` in the config (0.1.31, opt-in) runs such a pair on the first card alone when it |
| holds every profiled expert: with the R9700 first, 4K prompts 1,776 tok/s (split: 1,244) and decode ~60 tok/s |
| (split: ~51), the tokens of the R9700 alone (docs/MULTI_GPU.md). |
| - **Speed switches (engine 0.1.32, measured on the R9700 / 9070 XT with the Coder IQ1_M pack):** |
| - the MoE router (`router_top10`) runs a HIP kernel without its serial FP64 sum and block barriers by default: the |
| same ids and weights bit for bit (`hip_router_fast` checks 65,536 rows), 39 -> 9-12 us per call, the same greedy |
| tokens (5 + 5 starts). Decode: 62.4 -> 70.0 tok/s on the R9700 and +4% on the 9070 XT in one A/B here (ROCm 10.2 |
| nightly); a user's repeated A/B/A/B with setup's TheRock 7.10 wheels and an i5-12600K measured +1-4%, inside a |
| 12-20% run-to-run spread (#432) - how much it gains depends on how much of decode the router is on that setup. |
| `STRATA_HIP_ROUTER_OLD=1` runs the portable kernel. |
| - `STRATA_HIP_WMMA=1` (opt-in, gfx12 only, int8 KV): the prompt path's QSA attention on RDNA4 matrix cores |
| (`v_wmma_f32_16x16x16_f16`, FP16 hi + lo halves like the CUDA tensor-core kernel; `hip_prompt_attn_wmma` bounds it |
| against the FP32 kernel and FP64). 7.2-7.5x the portable kernel; R9700 prompts 4K 1,784 -> 2,427 tok/s, 16K |
| 1,797 -> 2,700 (a user with TheRock 7.10: +29-34%, #432). Not bitwise: greedy text differs from token ~50 on, as with the CUDA tensor-core attention. With it |
| the prompt path's expert ring is 96 slots (as STRATA_PREFILL_RING=96): with the default 384 the 9070 XT's 4K |
| prompts fell to 718 tok/s; with 96 they gain (1,017 -> 1,211; 16K 1,518 -> 2,032). PR #329 |
| (bsorensen110) contributed an equivalent gfx12 WMMA kernel of the same speed (within 1%); this one also masks KV |
| pages that KV streaming has not made resident, as the decode kernel does. |
| - The prompt path's QSA top-k picks its kernel by the blocks a query actually has, not the cache's capacity (#337, |
| bsorensen110): the same ids, on by default with AMD (NVIDIA keeps its capacity rule: there the 64K prompts read |
| 1-3% slower with it). `STRATA_SELECT_WMMA=1` (opt-in, gfx12) adds #337's |
| matrix-core block scorer; it selects slightly differently (254 of 256 queries the same) and gained +1.5% on 16K |
| prompts at a 262K context on the R9700. |
| - **Known:** rarely (about 1 start in 10) a HIP run's greedy output differs from another start's at some token, on |
| one card or two and on engine 0.1.29 as well; not yet explained. |
| - **Not validated:** images, long contexts beyond 16K, answer-quality benchmarks. |
| |
| ## Community-validated cards |
| |
| Run by their owners, not on the maintainers' machines; setup accepts them like gfx1100 / gfx1201. setup ships a |
| hipBLASLt table for gfx1200 (the numbers below); gfx1101 has none (make one with [Tuning table](#tuning-table) |
| and compare the prompt speed with and without it). |
| |
| - **gfx1101, RX 7800 XT 16 GB** (jhohertz, #254; engine 0.1.29, Ryzen 9 5950X, 121 GiB RAM, system ROCm with |
| hipBLASLt 1.4.1): `./setup.sh --backend hip` detected the card and compiled the engine; `strata-device --selftest` |
| passed; ctest 30/32 (`ple_parity` needs the Q2_0 fixture, `platform_memory_test` the memlock limit). Coder IQ1_M, |
| 64K context, MTP, with a table the owner calibrated: fresh prompts of 4K-9K tokens at 898-953 tok/s, decode |
| 38-44 tok/s (128 tokens). |
| - **gfx1200, RX 9060 XT 16 GB** (Efeisot, #256 after #176; engine 0.1.29, Ryzen 9 7950X, 64 GiB RAM, ROCm 7.2 with |
| hipBLASLt 1.2.2): ctest 32/32 (without `ple_parity` and `platform_memory_test`). Coder IQ1_M, greedy, MTP |
| `--spec 4 --spec-min-p 0.5`, with a table the owner calibrated: |
| |
| | prompt | prompt speed | decode | |
| |---|---|---| |
| | 2,374 tokens | 540 tok/s | 27.1 tok/s (0.69 draft acceptance; 31.0 at 1.00) | |
| | 65K (`--kv int8 --kv-resident 65536 --prefill 16384`) | 748-753 tok/s | 26-40 tok/s by acceptance | |
| | 130K (the same flags) | 725 tok/s | - | |
| |
| Greedy output was the same across runs. For comparison, llama.cpp's HIP build measured 20 tok/s decode and |
| 450 tok/s prompt on that card. |
| |
| Since 0.1.38 setup ships that table (`tools/hip/gfx1200-hipblaslt-100202.txt`). On engine 0.1.31, one binary (ctest |
| 37/37, without `ple_parity` and `platform_memory_test`), prompts of repeated code blocks, prefill 2560 for |
| the short prompt and 16384 with `--kv int8 --kv-resident 65536` otherwise, two runs each - prompt speed |
| with the table vs plain hipBLAS (no table): |
| |
| | prompt | with the table | plain hipBLAS | gain | |
| |---|---|---|---| |
| | 2,374 tokens | 597-599 tok/s | 396-398 tok/s | 1.50x | |
| | 65,045 tokens | 754-757 tok/s | 380 tok/s | 1.99x | |
| | 130,091 tokens (the same flags) | 719 tok/s | 363-370 tok/s | 1.98x | |
| |
| ## RDNA2 (gfx1030) |
| |
| The RX 6800 / 6800 XT / 6900 XT / 6950 XT run the same kernels: wave32 and 64 KiB of LDS per workgroup. The one |
| difference is the dot instruction: RDNA2 has no `v_dot4_i32_iu8` (gfx11 and newer), so |
| `dp4a` uses the plain signed `v_dot4_i32_i8` through `__builtin_amdgcn_sdot4`, which compiles to a single |
| `v_dot4c_i32_i8` (same signed x signed byte products, modulo 2^32). There is no WMMA; the QSA scorer takes the |
| same ordered FP32 fallback as gfx1100. CMake lists gfx1030 as unvalidated (the build warns) until a maintainer has |
| run it; the report below is from a community machine: an RX 6900 XT 16 GB (gfx1030), an i7-13700KF (8 P-cores and |
| 8 E-cores, AVX2, no AVX-512), 63 GB RAM, NixOS, ROCm 7.2.3 from nixpkgs (clang 22, hipBLAS 3.2). |
|
|
| - **Build and tests** (engine 0.1.26): a complete HIP build for gfx1030, made by hand with cmake and ROCm's own |
| `clang++` (the nixpkgs ROCm is not an `/opt/rocm` tree, so setup's `build_engine_hip` was not exercised; the |
| binary was placed in `engine/` for setup to use). All 28 registered ctest tests pass on the card, including |
| `hip_device_selftest`. An engine 0.1.30 build of this branch configures, builds and runs clean on the same card; |
| the speeds below are measured on it. |
| - **End to end** (Swift 1.5 IQ3_XXS, `--context 131072` with `--kv-resident 32768 --adapt-every 1 |
| --vram-reserve-mib 1024`, 200 greedy tokens): 38-42 tok/s decode with the default 15 CPU pool workers (one per |
| physical core except the host thread; 8 workers: 36; 24 = every logical core: 27), consistent even with the |
| 131,072-token context full. |
| - **Prefill** (the same 15-worker configuration): 246 tok/s on a 2,000-token prompt, 330-339 tok/s at the auto |
| 8,192-token chunk (7,997 and 15,967 tokens; time to first token 8.9 and 48.3 s), 133 tok/s on a 522-token |
| prompt - short prompts are fixed overhead (19-25 tok/s on 34 tokens). Decode after a 16K prefill holds at |
| 45.5 tok/s. |
| - **16 GB card:** the expert cache holds 6,310 slots (10.2 GiB) at 16K context and 5,247 slots at 131,072, where |
| setup keeps the KV cache in VRAM because IQ3_XXS needs about 60 GB of RAM plus the cache to stream it. The |
| `--pcie-frac 0` and `--adapt-every 0` of the 7900 XTX configuration in |
| [AMD_HIP_PERFORMANCE.md](AMD_HIP_PERFORMANCE.md) cost 8.6 and 13.6 tok/s here (30 with the defaults): keep the |
| defaults on a 16 GB card. |
| - **hipBLASLt:** ROCm's hipBLASLt ships no gfx1030 kernels, so there is no table and the plain hipBLAS path runs. |
| - **Not validated:** gfx1031 / gfx1032 (the same `dp4a` path, no hardware report), setup's own build path and the |
| `gfx103X-all` wheels on gfx1030, images, answer-quality benchmarks. |
|
|
| ## Tuning table |
|
|
| A hipBLASLt table holds solution ids that are valid only for one GPU architecture and one hipBLASLt version, so it |
| is calibrated on the card with the ROCm the engine runs with. The shipped gfx1100 table's rows are the engine's |
| dense GEMM shapes; to calibrate them for another card or version (a few minutes): |
|
|
| ```sh |
| cmake --build build-hip --target tune_hipblaslt |
| CASES=$(awk 'NR>2 {printf " --case %s,%s,%s,%s,%s", $1, $5, $2, $3, $4}' tools/hip/gfx1100-hipblaslt-100200.txt) |
| ./build-hip/tune_hipblaslt $CASES --tuning-out table.txt |
| ``` |
|
|
| The file's second line names the architecture and version (`STRATA_HIPBLASLT_TUNING_V1 gfx1201 100401`); save it |
| as `tools/hip/<arch>-hipblaslt-<version>.txt` for setup, or point `STRATA_HIPBLASLT_TUNING` at it. Compare the |
| prompt speed with and without it before keeping it. |
|
|
| Shipped tables: |
|
|
| - `gfx1100-hipblaslt-100100.txt`, `gfx1100-hipblaslt-100200.txt`: RX 7900 XTX. |
| - `gfx1201-hipblaslt-100500.txt`: Radeon AI PRO R9700 (gfx1201, 32 GB), calibrated with ROCm 10.2.0a20260914 |
| (AMD's `gfx120X-all` nightly, hipBLASLt 1.5.0, library build `d3164197`). 16 dense GEMM geometries at T=4096 and |
| T=8192, 32 rows. setup uses it only when the installed hipBLASLt reports 1.5.0 (it is found in `/opt/rocm` |
| when that is a system ROCm 7 or newer). The version number is the only thing the engine can check, so another |
| 1.5.0 build could number its solutions differently. `hip_prefill_hipblaslt_gemm` (with `STRATA_HIPBLASLT_TUNING` |
| set) is a smoke test: it refuses a table for another architecture or version and checks that one BF16 and one |
| F16 row exist and agree with hipBLASEx, which covers 2 of the 32 rows. It does not prove that a solution id is |
| valid: the engine falls back to hipBLASEx for an id the library rejects, and the test still passes. Run it with |
| `STRATA_HIPBLASLT_VERBOSE=1` and look for `fallbacks=0` in its summary line, and recalibrate with |
| `tune_hipblaslt` before using this table with a different 1.5.0 build. |
|
|
| ## Original backend validation (PR #94) |
|
|
| The following is historical validation of the original backend, not a fresh |
| test count for this replacement. Current build/test evidence and performance |
| limits are recorded in [AMD_HIP_PERFORMANCE.md](AMD_HIP_PERFORMANCE.md). |
|
|
| Tested against upstream `c1e903310f211e6630780c3bd2038778c071c68d` (0.1.20), |
| with pinned llama.cpp `3cf03257f219afbe7334045ff7c6a06ac68c627d`. |
| The host has an RX 7900 XTX (24 GiB), Ryzen 9 7950X3D, 64 GiB system RAM, |
| and Fedora-family Linux with ROCm 7.1.52802 / Clang 20. The original model used |
| for the smoke check was on mechanical RAID0; cold page faults were slow. Keep |
| latency-sensitive data on SSD and measure warm residency separately. |
|
|
| - Complete HIP Release build, including the executable and parity targets. |
| - All 25 tests selected by the command below passed on 2026-09-29. |
| - Intrinsic parity includes signed packed-byte dot products and overflow. |
| - CPU/GPU mapped-memory handoff covers delayed publication, changing payloads, |
| and repeated graph replay in both directions. |
| - GR coverage includes the maximum eight-token fused batch and graph replay |
| after inputs change. Native QSA checks the scalar HIP fallback. |
| - Q8_K quantization is byte-exact against the existing reference. HIP's |
| `__fmul_rn` can become ordinary multiplication; disabling contraction in |
| `quantize_act.cu` preserves the separately rounded product before the magic |
| rounding bias is added. CUDA's compile flags remain unchanged. |
| - The mmap regression covers canonical and differing native layer sizes, |
| truncated/oversized files, lookup bounds, and reopening after errors. |
| - CUDA 13 / GCC 15 / sm_89 executable build passed; GR, sampler, quantization, |
| and mmap tests passed on the same tree. This is a regression check, not a |
| claim of complete CUDA inference validation. |
| - CPU-only `strata-plan` build passed without either GPU backend. |
| - Real GSQ-RCO IQ3_XXS server smoke with mmap, the shipped expert profile, |
| automatic GPU cache, and MTP passed arithmetic, executable Python addition, |
| system-marker recall, and 1,170-token batched-prefill recall. All four requests |
| ended normally with correct answers; one engine process served them all. |
| The cache held 11,142 experts (18.08 GiB), with 874 MiB of VRAM free after |
| startup. These short checks do not establish a throughput or quality benchmark. |
| |
| Build all targets before running the registered focused checks: |
| |
| ```sh |
| cmake --build build-hip -j2 |
| ctest --test-dir build-hip --output-on-failure --timeout 60 \ |
| -E '^(ple_parity|platform_memory_test)$' |
| ``` |
| |
| `ple_parity` requires an external model fixture. `platform_memory_test` requests |
| 256 MiB of locked memory; the test shell's 8 MiB memlock limit prevented that |
| test from passing. These are explicit exclusions, not skipped tests counted as |
| passes. Additional unpublished upstream fixture suites are not covered. |
| |
| The published performance table explicitly identifies the measured development |
| snapshot; it must not be read as a benchmark of every subsequent rebase. Vision, long-context stress, broad answer-quality |
| equivalence, other AMD cards, and mixed-vendor inference are not validated here. |
| |
