Skip to content

Run Edge0-145B Q3_K on AMD HIP (Linux, gfx1100) - #1

Draft
ihavespoons wants to merge 7 commits into
mainfrom
hip-latent-q3k
Draft

ihavespoons wants to merge 7 commits into
mainfrom
hip-latent-q3k

Conversation

@ihavespoons

@ihavespoons ihavespoons commented Nov 13, 2026 •

Copy link
Copy Markdown

What this is

This branch carries the AMD work for Edge0-145B Q3_K on Linux with Radeon RX 7900 XTX cards (gfx1100). The pull request stays on this fork while the work continues.

Contents

  • 78fe2d4 is feat(strata): build and run edge0-145b on AMD HIP (Windows, gfx1100) Edge0-AI/Edge0#196 by its original author, unchanged. It makes the latent-MoE sources build and run on HIP, with the fused latent kernels off.
  • ea78b76 and 15fa560 add "Build (Linux, AMD, experimental)" to edge0-strata/EDGE0.md: the build command, the tuning table step, the server configuration, and the measured figures.
  • 17c02cd adds src/prefill/lat_dense_w11.cu: the D and U dense projections of a latent layer on the gfx11 matrix cores (v_wmma_i32_16x16x16_iu8). prefill.cpp calls it on HIP in place of iq_mmvq and latent_u_grouped.
  • 4cb5224 adds src/prefill/lat_q3k_w11.cu: the expert stage of a latent layer on the same matrix cores. Gate/up reads Q3_K, applies SwiGLU and rounds H to q8_1. Down reads IQ4_NL. STRATA_LATENT_FUSED=1 with STRATA_LAT_FUSED_Q3K=1 selects it, as on CUDA.
  • 724ba73 updates the Linux section for the expert stage and the two-card split.
  • 49d5def fixes the layer split on HIP. The latent D/U pair was one process-wide pointer on the first device, and STRATA_EMBED_VRAM=1 put the token embedding table there too. A later stage and the draft layer were given those pointers, and a kernel cannot read another device's VRAM. Each device now has its own D/U pair for its layers, and on a split the embedding table stays in mapped host memory.

Measured on Linux

Ryzen 5 5600X, 64 GB RAM, ROCm 7.2.4 in rocm/dev-ubuntu-24.04:7.2.4-complete, one RX 7900 XTX on PCIe 4.0 x16, --vram-reserve-mib 700, 6652 experts resident. Thinking off, one run each.

Configuration Chat decode tok/s 5.3K prompt tok/s 21K prompt tok/s
Edge0-AI#196 as it is (STRATA_LATENT_FUSED_DU=0, no tuning table) 47.7 to 57.1 214.3 216.7
plus STRATA_LATENT_MMQ=1 46.4 to 55.1 269.1 280.2
plus the gfx11 D and U kernels not measured 724.9 not measured
plus the hipBLASLt tuning table 46.2 to 51.1 1202.0 1316.1
plus the gfx11 expert stage in place of MMQ not measured 1181.8 1302.1

The gfx11 expert stage and the MMQ stage ran in different sessions. The same MMQ configuration measured 1157.2 and 1202.0 tok/s on the 5.3K prompt, so the two stages are level within the run-to-run spread. The expert stage takes 778 ms of a 4002 ms GPU timeline.

The same card on a PCIe 3.0 x1 link decodes at 17.8 to 23.3 tok/s.

Two cards in a layer split with --resident-experts, layers 0 to 21 on the x1 card and 22 to 47 on the x16 card: 27.0 GiB of the 37.3 GiB of experts in VRAM, 21.4 GiB kept in RAM, 32.5 GB of RAM for the process. Chat answers decode at 32.2 to 51.2 tok/s. A 5.3K prompt runs at 272.9 to 284.6 tok/s and a 21K prompt at 413.2 tok/s.

Checks of the new kernels

  • Layer 0 of a 5268-token prompt, against the earlier kernels: 99.9987% of the D stage's q8_1 codes are equal and the rest differ by 1. The dumped U rows agree to 2e-7.
  • First-token top-8 log-probabilities on four prompts move by at most 0.57. The two existing expert stages (STRATA_LATENT_MMQ set and unset) differ from each other by up to 2.41 on the same prompts.
  • The gfx11 expert stage moves the same log-probabilities by at most 0.48 against the MMQ stage.
  • 40 short questions with one exact answer: 39 of 40 correct.

Route flags

STRATA_ROUTE_TAIL_SKIP=7 with STRATA_ROUTE_PRIOR=1.5 changes the output by design. On 300-token answers the decode speed went from 51.3 to 53.4 tok/s to 56.2 to 59.8 tok/s. The same 40 questions scored 40 of 40. Forty short questions cannot show a small quality loss.

Found

  • The Linux HIP build needs -DSTRATA_PREFILL_MMQ=ON. Without it the link stops with undefined symbols such as strata::prefill::fused_q3k::experts_latent_q3k.
  • A build started by hand has no tuning/ folder beside the engine, so the dense prompt products run on plain hipBLAS.
  • STRATA_LAT_MOE_DUMP's a-yl.bin is not stable: two runs of the same build and prompt differ in it (correlation 0.85) while their d-xlq.bin and their replies are identical.
  • Before 49d5def, a layer split failed in several ways: verify: timed out at layer 21 or 22, Memory access fault by GPU node-2, an engine exit with code -9 during prompts, and a main thread that waited in hipStreamSynchronize in MtpDrafter::prefill. The kernel log showed sq_intr: error on the reading card each time. A test program gives the cause: a kernel on one device that reads a buffer of the other device gets "an illegal memory access was encountered".
  • A blocking hipMemset of another device's buffer returns before that device sees the bytes. The engine has no such call, and the compat layer now runs these calls on the buffer's own device.
  • --pcie-mode dma on a split was not tested again after the fix.

Not done

  • The gfx11 kernels reach about a quarter of the card's integer matrix rate by my estimate. I did not profile inside them.
  • The HIP unit tests and the latent_*_verify tools were not run.
  • Edge0 120B, the NVFP4 folders and the Q4_0 projection file were not tried.

🤖 Generated with Claude Code

ppkrol and others added 6 commits November 12, 2026 19:54
edge0-strata keeps upstream Strata's HIP backend, but the latent-MoE
additions did not build or run with it.

Build (HIP only):
- CMakeLists.txt: compile lat_q3k_*.cu and moe_fused_q3k.cu in the HIP MMQ
  branch (their entry points are stubs there) and add latent_dense.cpp and
  latent_moe.cpp to strata_engine for HIP
- moe_fused.cu, moe_fused_q3k.cu: the latent kernel launches were outside
  the !__HIPCC__ guard that holds the kernels
- verify.cpp, prefill.cpp: two debug dumps use cudaGraph* / cudaPointer*
  names the HIP shim does not map

Run (HIP only):
- verify.cpp: the D stage's two D2H cudaMemcpyAsync calls and the
  "verify top3" dump (a synchronous cudaMemcpy) ran inside the verify
  window's graph capture, which HIP rejects ("end capture: operation
  failed due to a previous error during capture"); on HIP the copies go
  through copy_bytes_from_mapped and the dump is skipped

Docs: EDGE0.md describes the build and the configuration.

The fused latent kernels stay CUDA-only; on HIP the model runs with
STRATA_LATENT_FUSED=0, STRATA_LATENT_FUSED_DU=0, STRATA_LAT_FUSED_Q3K=0.
Tested on one RX 7900 XTX under Windows 11. The CUDA build was not rebuilt.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VmeXwKqR7k1jsojD2bUxve
The HIP build needs -DSTRATA_PREFILL_MMQ=ON on Linux. Without it the
link stops with undefined symbols from the latent prefill sources.

Add the Linux steps and the figures measured on one RX 7900 XTX with
ROCm 7.2.4. Record two limits: a layer split across two cards stops
with a GPU memory access fault, and a card on a PCIe x1 link decodes
at less than half the speed of the same card on x16.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
On HIP the prompt path ran the two dense projections of a latent layer
on single-row kernels: iq_mmvq for D and latent_u_grouped for U. They
hold 82.6M of the 117M multiply-adds per token and layer, and they took
13.0 s of a 25.7 s GPU timeline for a 5268-token prompt.

lat_dense_w11.cu runs both as Q8_0 x q8_1 integer dots on
v_wmma_i32_16x16x16_iu8, 128 weight rows against a 64-row tile per work
item. D quantizes its products to the xl rows' q8_1 blocks as
q8_1_store does, and writes the FP32 rows too when STRATA_LATENT_MMQ=1
needs them. STRATA_LATENT_FUSED_DU=0 keeps the earlier kernels. The
Q4_0 projection file keeps them too.

Measured on one RX 7900 XTX (gfx1100, PCIe 4.0 x16), ROCm 7.2.4:

- 5268-token prompt, native A stage: 213.4 -> 403.0 tok/s
- 5268-token prompt, STRATA_LATENT_MMQ=1: 266.2 -> 682.3 tok/s
- 19494-token prompt, STRATA_LATENT_MMQ=1: 280.4 -> 780.4 tok/s

Checked against the earlier kernels on the same prompt: 99.9987% of the
D stage's q8_1 codes are equal and the rest differ by 1; the dumped U
rows agree to 2e-7. The first token's top-8 log-probabilities move by
at most 0.67 on four prompts, and the two earlier A stages differ from
each other by up to 2.41 on the same prompts.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
… Linux

A build started by hand has no tuning/ folder beside the engine, so the
dense prompt products ran on plain hipBLAS. With the table for
hipBLASLt 1.2.2 a 5263-token prompt runs at 1150 tok/s, against 725
without it.

Update the measured figures for the gfx11 D and U kernels.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
latent_kernel_q3k needs mma.sync, so HIP had no fused expert path for the
q3k pack. lat_q3k_w11.cu adds one for gfx11: gate/up (Q3_K) with SwiGLU and
the q8_1 rounding of H, then down (IQ4_NL), each expert read where it is.

A Q3_K sub-block has one 6-bit scale. The kernel multiplies it into the
int8 operand as (4 - code) * (scale - 32), which is in [-128, 124], so the
matrix core sums a whole 32-value activation block exactly.

STRATA_LATENT_FUSED=1 STRATA_LAT_FUSED_Q3K=1 selects it. On HIP the
lat_q3k_fused path (CUDA-only stubs) is now off.

RX 7900 XTX, 5263-token prompt: 1181.8 tok/s against 1157.2 on the MMQ
A stage. First-token log-probabilities differ from the MMQ path by at
most 0.48 over four prompts.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
… and embeddings

A kernel cannot read another device's VRAM. HIP answers "an illegal memory
access was encountered", the kernel log shows sq_intr errors on the reading
card, and the graph never finishes. Two tables were on the first device
only, and a later stage was given their pointers:

- The latent D/U pair was one process-wide pointer (latent_shared_d). It is
  now one pair per device. A later stage uploads the pair of the layers it
  runs before its expert cache is sized. With --resident-experts the window
  stopped with "verify: timed out at layer 22".
- STRATA_EMBED_VRAM=1 put the token embedding table in the first device's
  VRAM, and the draft layer gathers its rows on the last device. On a layer
  split the table now stays in mapped host memory. The main thread waited
  in hipStreamSynchronize in MtpDrafter::prefill.

The HIP compat layer now runs a blocking memcpy or memset on the buffer's
own device. A blocking hipMemset of another device's buffer returns before
that device sees the bytes. The engine has no such call today
(STRATA_HIP_XDEV_LOG=1 printed none), so this guards later code.

Two RX 7900 XTX, --resident-experts: 27.0 GiB of experts in VRAM, 21.4 GiB
kept in RAM where one card keeps 37.3 GiB, chat decode 32.2 to 51.2 tok/s.

Co-Authored-By: Claude Fable 5.2 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants