Repository navigation
Run Edge0-145B Q3_K on AMD HIP (Linux, gfx1100) - #1
Draft
ihavespoons wants to merge 7 commits into
Draft
ihavespoons wants to merge 7 commits into
ihavespoons wants to merge 7 commits into
Conversation
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>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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
78fe2d4is 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.ea78b76and15fa560add "Build (Linux, AMD, experimental)" toedge0-strata/EDGE0.md: the build command, the tuning table step, the server configuration, and the measured figures.17c02cdaddssrc/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.cppcalls it on HIP in place ofiq_mmvqandlatent_u_grouped.4cb5224addssrc/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=1withSTRATA_LAT_FUSED_Q3K=1selects it, as on CUDA.724ba73updates the Linux section for the expert stage and the two-card split.49d5deffixes the layer split on HIP. The latent D/U pair was one process-wide pointer on the first device, andSTRATA_EMBED_VRAM=1put 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.STRATA_LATENT_FUSED_DU=0, no tuning table)STRATA_LATENT_MMQ=1The 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
STRATA_LATENT_MMQset and unset) differ from each other by up to 2.41 on the same prompts.Route flags
STRATA_ROUTE_TAIL_SKIP=7withSTRATA_ROUTE_PRIOR=1.5changes 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
-DSTRATA_PREFILL_MMQ=ON. Without it the link stops with undefined symbols such asstrata::prefill::fused_q3k::experts_latent_q3k.tuning/folder beside the engine, so the dense prompt products run on plain hipBLAS.STRATA_LAT_MOE_DUMP'sa-yl.binis not stable: two runs of the same build and prompt differ in it (correlation 0.85) while theird-xlq.binand their replies are identical.49d5def, a layer split failed in several ways:verify: timed out at layer 21or22,Memory access fault by GPU node-2, an engine exit with code -9 during prompts, and a main thread that waited inhipStreamSynchronizeinMtpDrafter::prefill. The kernel log showedsq_intr: erroron 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".hipMemsetof 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 dmaon a split was not tested again after the fix.Not done
latent_*_verifytools were not run.🤖 Generated with Claude Code