Support for AMD GPUs (Vulkan or ROCm/HIP) for Edge0 120B / 145B #197

Activity

1. Thomas12323

   added

   enhancementNew feature or requestNew feature or request

   on Nov 12, 2026
2. shuk777 commented on Nov 13, 2026

   

   shuk777

   on Nov 13, 2026

   Collaborator

   More actions

   Thanks for your interest in our work! Our next goal is of course supporting more environments, thanks for offering to help test. There is already a pr #196 focusing on AMD support, our team will review it immediately. Will let you know once AMD support is ready!
3. gangadharjannu commented on Nov 13, 2026

   

   gangadharjannu

   on Nov 13, 2026

   More actions

   Mine is R9700 AI PRO card, I am also happy to test
4. JuggyBob commented on Nov 13, 2026

   

   JuggyBob

   on Nov 13, 2026

   More actions

   Adding a second AMD data point, since PR #196 (#196) covers gfx1100 with 24 GB VRAM and this is a 16 GB card.

   Machine (runs upstream Strata's HIP build today):

   • GPU: Radeon RX 7600 XT 16 GB — gfx1102, RDNA3, 32 CUs, 288 GB/s
   • CPU: Ryzen 5 7600X (6C/12T)
   • RAM: 31.1 GB
   • OS: Windows 11 (10.0.26300), driver 32.0.32015.7029

   For context on what gfx1102 already does here: this box serves Qwen3.8-Flash-Next Coder IQ1_M (23.4 GiB of expert weights) on upstream Strata's hip-rdna3 engine build at 39.5–44.4 tok/s decode and 320–341 tok/s prefill at 256K context, so the RDNA3 wave32 path is working, not just compiling.

   Questions for the AMD branch:

   1. Does STRATA_HIP_ARCHS=gfx1102 build cleanly, or is anything in PR feat(strata): build and run edge0-145b on AMD HIP (Windows, gfx1100) #196 gfx1100-specific? Same generation, but the CU count and L2 differ.
   2. Memory sizing. Edge0-145B-q3k is 40.81 GiB of weights; 16 GB VRAM + 31 GB RAM leaves no headroom at all once Windows and the desktop are up. Is that expected to run via --mmap-experts / --resident-budget-gib, or is 145B effectively out of scope below ~24 GB VRAM? Edge0-120B-q3k (27.98 GiB) fits this machine's RAM comfortably by comparison, so it may be the realistic target here.
   3. RDNA3 has no FP4 MMA, so Q3_K is the only format for us. Is the Q3_K latent path the one the HIP branch actually exercises (STRATA_LATENT_FUSED=0, STRATA_LAT_FUSED_Q3K=0 as in the PR)?

   What I can contribute if useful: a gfx1102 build, tools/strata_bench.py numbers for 145B-q3k and 120B-q3k at 16 GB VRAM / 31 GB RAM, and a sweep of the residency settings at that budget. I have no NVIDIA card, so I cannot provide a CUDA comparison.

   I have not built or run anything yet — the numbers above come from upstream Strata, not from Edge0.
5. GravLabs commented on Nov 15, 2026

   

   GravLabs

   on Nov 15, 2026

   More actions

   RDNA4 (gfx1201) data point + a branch: edge0-145b Q3_K on 2x Radeon AI PRO R9700, with the latent D/U projections on the gfx12 matrix cores

   Branch, offered as-is (use any of it, or none): https://github.com/GravLabs/Edge0/tree/amd-gfx12-rdna4

   It sits on top of the AMD HIP work from #196 (78fe2d4) and the gfx11 latent kernels from diffsec/Edge0 (2484645, by @ihavespoons). That work isn't in main yet, so this is a branch, not a PR. Our four commits on top:

   • 9dce143 strata generate: refuse --peer-device on a latent pack. This is a bug we hit and reproduced (details below).
   • ec21935 cmake: link strata_kernels_cpu into the ten tools that need lat_gu_colscale_host. They failed to link on HIP.
   • f77909b prefill: the latent D/U projections on gfx1200/gfx1201. lat_dense_w11.cu now also builds for the gfx12 int8 WMMA (v_wmma_i32_16x16x16_iu8, gfx12 layout), plus a HIP unit test, hip_lat_dense.
   • 9940228 hip_lat_dense: guard bands; line wraps only in the kernel.

   Environment

   Threadripper PRO 9955WX, 188 GB RAM, 2x Radeon AI PRO R9700 (gfx1201, 32 GB, 300 W cap), ROCm 7.2.3, Ubuntu 24.04. Model Edge0-145B-q3k, STRATA_LATENT_DENSE_FMT=8, 262K context. The two cards run as a layer split (gpu: [0, 1], --resident-experts, STRATA_VERIFY_ALL_RESIDENT=0, --vram-reserve-later-mib 1322), with STRATA_LATENT_MMQ=1.

   Bug: --peer-device on a latent pack returns no token (reproduced)

   With --peer-device 1 on 2x R9700, every request read the prompt, hit non-finite logits (#879), retried once and failed with 0 tokens. That held for every request: 10 of 10 in the service journal, plus all of our test runs, at prompt lengths from 54 to 19K tokens.

   • Same config on one card without the flag: coherent output. The last residual row at prompt end has max |x| ~4, against ~1.5e7 with the flag.
   • Cause, from reading the source (prefill.cpp, peer_experts.cpp):
     • Prompt path: the peer is assigned its experts' rows, but the latent branch never launches the peer, so those rows are never computed.
     • Decode path: PeerExperts::launch runs the expert on the full-width row without the latent D projection, labels or U projection.
   • Not HIP-specific by reading, but we did not run it on CUDA.
   • The guard exits 2 with a fact/why/fix message, verified on the GPUs. A latent-aware peer tier would be the real fix.

   The layer split on the #196 commit alone also crashed here (verify: layer 24 never rang (an illegal memory access)). It needs diffsec's 49d5def (one latent D/U pair per card), which this branch includes. --pcie-mode dma on the split was flaky (verify timeouts), so we didn't use it.

   Results (same box, same config except STRATA_LATENT_FUSED_DU=1; 2 interleaved runs each, no profiler)

   prompt tok/s	generic path	gfx12 D/U	speed-up
   ~1K	248	655	2.6x
   ~16K	352	2,073	5.9x
   ~64K	467	2,620	5.6x
   decode, single stream	39-40	39-41	~1.0x

   The two runs agreed to within 0.1%. Profile of the generic path at 16K (rocprofv3): latent U 59%, latent D (mmvq_kernel<8>) 31%, A stage 3%, hipBLASLt 2% of prompt GPU time. Decode is not D/U-bound: 77% of decode GPU time is wait_flag_ge_kernel, the cross-card wait in the layer split.

   Same gfx12 D/U config at longer contexts and under concurrency. Server parallel is marked per row; every run returned tokens.

   test	result
   prompt ~128K (134K tok), parallel 8	2,430 tok/s (TTFT 55 s)
   prompt ~183.5K (188K tok), parallel 8	2,410 tok/s (TTFT 78 s)
   decode at depth 0 / 64K / 128K, parallel 8	41.7 / 41.9 / 43.7 tok/s
   2K prompts, concurrency 1 / 2 / 4 / 8 / 16, total output, parallel 8	34.5 / 48.7 / 61.5 / 80.5 / 81.7 tok/s
   8 concurrent ~129K prompts, parallel 4	2,712 tok/s aggregate; TTFT 48 / 215 / 382 s (min / median / max)
   8 concurrent ~129K prompts, parallel 8	2,390 tok/s aggregate; TTFT 56 / 245 / 433 s

   What we read from it:

   • Prompt: a prompt costs about 1.2 s fixed plus L / ~2,300-2,430 tok/s, flat from 16K to 188K with no visible attention penalty. The fixed part is why short prompts are slow.
   • Decode: flat with depth.
   • Concurrency: prompt reads run one session at a time. The 8-session aggregate is about the single-stream rate, so the n-th session waits about n x 48-54 s for its first token. Total decode scales 2.3x from 1 to 8 sessions.
   • parallel: 4 beats 8 here, by 13.5% on the 8 x 128K run. Single-stream prompts at parallel 8 were also slower than the parallel 4 table above, e.g. 500 vs 655 tok/s at 1K.

   For scale: the same box running Radiance Q4

   This box's everyday engine is Radiance 1.17.0 serving Qwen3.8-Flash-Next at 4 bits (fp8-iq4r-moe), tensor-parallel across both cards. The model and engine are different, so treat this as context, not a like-for-like comparison. Both were measured with the same benchmark tools on the same 2x R9700, on 2026-11-14/15.

   test	Edge0-145B, gfx12 D/U	Radiance Q4
   prompt ~1K / 16K / 64K tok/s	655 / 2,073 / 2,620 (parallel 4)	6,553 / 7,946 / 7,755
   prompt ~128K tok/s (TTFT)	2,430 (55 s)	7,595 (17.7 s)
   prompt ~183.5K tok/s (TTFT)	2,410 (78 s)	7,497 (25.1 s)
   decode at depth 0 / 64K / 128K tok/s	41.7 / 41.9 / 43.7	139.5 / 130.8 / 135.4
   2K prompts, concurrency 1 / 2 / 4 / 8 / 16, total output tok/s	34.5 / 48.7 / 61.5 / 80.5 / 81.7	127.1 / 198.6 / 306.0 / 426.1 / 433.8
   per-request decode at concurrency 8 / 16	13.8 / 11.7	63.2 / 58.1
   TTFT p50 at concurrency 8 / 16 (2K prompts)	10.4 / 38.3 s	1.25 / 6.91 s
   long contexts, several at once: aggregate prompt tok/s	2,712 (8 x ~129K, parallel 4)	7,333 (4 x ~184K)
   long contexts, several at once: TTFT range	48-382 s	36-100 s

   Edge0 rows are at parallel 8 unless marked. Radiance's multi-session long-context row used 4 sessions at ~184K, not 8 at ~129K, so that row is the closest available shape, not an exact match.

   On the Edge0 side, the main remaining costs are prompt reads that serialize across sessions and the layer split's cross-card wait in decode.

   Validation

   Passed:

   • hip_lat_dense on both R9700s with two seeds: 18/18 cases at 0 ulp against a CPU reference of the kernel's terms. It reports SKIP rather than passing on a fallback.
   • gfx1100 code object of lat_dense_w11.cu: .text and .rodata byte-identical to the parent commit.
   • gfx1201 ISA: 8 v_wmma_i32_16x16x16_iu8 per instantiation, no s_trap, no scratch, ~99 VGPRs. Profiler traces show the kernels dispatching in the engine.
   • Engine first-token check: top-1 token and 16-token reply identical to the generic path in 21/21 pairs.
     • On ~1K, where fresh engines are bitwise deterministic, the port's own effect is a first-token KL of 7.4e-6 (top-5 tail tokens shift by up to 0.155 nats). This is expected, because the fused D quantizes straight to q8_1.
     • On ~3K, the generic path differs from itself by up to 0.56 nats (KL 6.8e-5) between runs, and every generic-vs-port pair falls inside that.
   • --peer-device refusal: exit 2 with the 3-part message.

   Failed / open:

   • On a 54-token prompt that doesn't reach the fused prefill (it goes through the verify/decode path), parallel 8 changed the 16th reply token compared with the parallel 4 reference ("…friendly. No" instead of "…friendly. Simple"). The top-1 first token was the same.
   • The same binary and config at parallel 4 reproduced the reference exactly. So this looks like batch-size-dependent decode numerics, not the port, but we haven't verified it further: the server returns log-probs only for the first token, so we couldn't see the margin at position 16.

   Not run:

   • gfx11 hardware: code-object identity only; we have no gfx11 card.
   • CUDA.
   • Windows.
   • Multi-token or teacher-forced KL.
   • The real-context and toolbench suites.
   • Occupancy tuning for gfx12 (STRATA_PF_OCC): the inherited gfx11 rule is unmeasured.
   • The A-stage (lat_q3k_w11.cu) on gfx12: not ported, because with STRATA_LATENT_MMQ=1 it is ~3% of prompt time.

   The work was AI-assisted (agents investigated, implemented and measured it). I reviewed it, and the numbers above are from actual runs. Happy to run something specific on this hardware if it helps.
6. balor1eye commented on Nov 15, 2026

   

   balor1eye

   on Nov 15, 2026

   More actions

   R9700 (gfx1201, 32 GB) data point on Linux: Edge0-145B-q3k runs with #196, and with gfx12 versions of the fused latent
   kernels on top (branch rebased onto current main, details: #196 (comment)) it reaches 1439 tok/s prompt at 4K,
   1718 at 16K and 59-68 tok/s decode (tools/strata_bench.py, thinking off). Strata's standard protocol passes 12/12 sanity
   checks, and our own 8-task agent suite passed 16/16 over two runs.

   Two things that cost me time on this card:

   • Ubuntu's ROCm 7.1 ships hipBLASLt 1.1, and the bundled gfx1201 tables are for 100202/100500, so prompt GEMMs silently
     fall back to plain hipBLAS. A table generated with tune_hipblaslt for 100100 is in the branch.
   • The R9700 is also my display GPU. With --vram-reserve-mib 2048 the expert cache took all but ~360 MB of VRAM and the
     KDE session crashed when the compositor rebuilt its surfaces. --vram-reserve-mib 6144 has been stable.

   @gangadharjannu you mentioned an R9700 earlier in this thread; if you try the branch, numbers from a second card would help.

   AI assistance: this work and comment were developed with Claude Code (Claude Opus 5.5); I reviewed it and ran the
   measurements on my own machine.