🍷 Winery Strata: fork of Niko1221/Strata (Winery Cuvée family, wine-red UI)

bbb6388 verified 26 days ago

31.8 kB

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) for the RX 7900 XT / XTX (RDNA3, gfx1100) and the RX 9070 / 9070 XT / Radeon AI PRO R9700 (RDNA4, gfx1201; see RDNA4). The RX 7800 XT / 7700 XT (gfx1101) and the RX 9060 XT (gfx1200) were validated by their owners (see 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). 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. 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.

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

./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). 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

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.

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; 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):

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, 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 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 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):

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.

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:

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.