rocm: Qwen3.6-35B-A3B GPU runtime — runs on Radeon 780M at ~14 tok/s, 5/5 top-1 parity#4
Open
marmyx77 wants to merge 10 commits into
Open
rocm: Qwen3.6-35B-A3B GPU runtime — runs on Radeon 780M at ~14 tok/s, 5/5 top-1 parity#4marmyx77 wants to merge 10 commits into
marmyx77 wants to merge 10 commits into
Conversation
added 10 commits
July 15, 2026 22:33
…warp) Without these the ROCm backend does not compile on ROCm 6.3.1 at all: - HIPBLAS_V2 + CUDA_R_* -> HIP_R_*: hipBLAS 2.x only declares the hipDataType-based GemmEx/GemmStridedBatchedEx overloads behind HIPBLAS_V2; the deprecated hipblasDatatype_t ones no longer match. - HIP_ENABLE_WARP_SYNC_BUILTINS: ROCm 6.3 gates __shfl_*_sync/__ballot_sync behind this macro. - __syncwarp(...) -> wavefront-scoped fence: no HIP equivalent exists; on RDNA3 wave32 lanes execute in lockstep so reconvergence is implicit and the fence supplies the compiler ordering the call relies on. Verified: ds4_rocm.o builds clean with hipcc 6.3.1, --offload-arch=gfx1100 (run on a Radeon 780M / gfx1103 via HSA_OVERRIDE_GFX_VERSION=11.0.0, since rocWMMA 1.6.0 rejects gfx1103).
ds4_gpu_host_memory_snapshot was a stub returning 0 on ROCm, which kept the adaptive expert-cache planner permanently unavailable there. Implement it for Linux: /proc/meminfo (MemTotal, MemFree, Inactive(file), Cached), /proc/self/status VmRSS, PSI /proc/pressure/memory for the pressure flags, and hipMemGetInfo total as the working-set envelope (on UMA APUs the GTT pool is the closest analogue of recommendedMaxWorkingSetSize). The Darwin field mapping is documented inline. Also store the planner's required floor and warn when the configured budget is below it (full fail-closed enforcement on ROCm still to be wired); the slab target stays a documented no-op (the ROCm cache is per-expert cudaMalloc, no slab allocator). Motivation, measured on a Radeon 780M / 64 GB / 48 GiB GTT running DeepSeek V4 Flash q2 with SSD streaming: with a fixed thin free floor the cache peaks after ~150-250 generated tokens and the kernel TTM evicts the GPU queue buffers themselves — [drm:amdgpu_cs_ioctl] *ERROR* Not enough memory for command submission! amdgpu: Freeing queue vital buffer ..., queue evicted — after which every HIP sync blocks forever with an idle GPU. A host-aware floor prevents it: 5.8-6.6 GiB floors froze or ran 2.97 tok/s; a 12 GiB floor ran 280-token soaks clean twice at 3.47 tok/s (best measured on that box; the wired cache competes with the page cache of the 81 GB GGUF, so the floor is a throughput optimum too, not just a safety line).
Any backend that can produce a live host-memory snapshot now gets the adaptive budget; backends without one keep the legacy fixed-fraction budget. Metal behavior is unchanged. This is the cross-backend step the planner entry in FORK_NOTES was waiting on: ROCm UMA APUs need the host-aware floor as much as Metal does (see the previous commit for the queue-eviction failure it prevents).
Multi-model range registry, per-value readback event ring with host-side waits (device-side hipStreamWaitEvent wedges HSA queues on this ROCm/APU stack), pinned-staging upload mutex, adaptive free-floor reserve, small-n prequantized q8 batch kernel, and the ds4_gpu_stream_expert_cache_stats zero stub required to link ds4-bench on non-Metal backends. Measured on Radeon 780M / 64GB / ROCm 6.3.1 (DeepSeek V4 Flash q2 SSD streaming): clean 280-token soaks at 3.47 tok/s with the adaptive floor; the event-ring/host-wait changes fix reproducible generation freezes.
Quoted -DDS4_BUILD_GIT_SHA values lose their quotes in recursive targets (strix-halo/cuda re-enter make with CFLAGS on the command line, adding a shell evaluation), breaking every ROCm/CUDA build with 'b60fbec8ddd7 undeclared'. A bare token survives any number of shell levels.
Port of metal/qwen35.metal to HIP (rocm/ds4_rocm_qwen.cuh): 19 device kernels + the 20 ds4_gpu_qwen35_* entry points. Metal simdgroups map to RDNA3 wave32 (HSA_OVERRIDE_GFX_VERSION=11.0.0); simd_sum/max/min become butterfly __shfl_xor reductions; byte-stride args mirror ds4_metal.m exactly. routed_moe_top8 delegates to the existing ROCm Q4_K MoE core (ds4_gpu_routed_moe_one_tensor) as a resident single top-8 pass. tests/test_qwen35_rocm.c validates every non-weight op against the scalar reference ds4_qwen_ref.c: split_q_gate, sigmoid_mul (elements/rows), rope, router (exact top-8 + weights), gqa decode/prefill, gated_delta_step (kd=128 fast path and kd=64 generic), gated_delta_sequence_128. All pass on the Radeon 780M — the core GDN recurrence matches to ~1e-8 relative. Weight-reading ops (conv, rmsnorm_gated, controls, embedding) still need a model-map-backed test and host wiring (Fase 3). + tests/qwen/normalize_qwen36_gguf.py: reproduces the blessed artifact (Q6_K ffn_down_exps -> Q4_K, output.weight -> Q8_0, pad id + chat template) via ggml reference quantizers.
conv_step, conv_sequence, rmsnorm_gated, gated_delta_controls, and dequant_embedding_q8_0 now validated against ds4_qwen_ref.c by registering a host weight buffer as the model map (UMA maps it device-side). All 22 per-op checks pass on the Radeon 780M; every kernel matches the scalar reference to <=1e-7 relative (GDN recurrence to ~1e-8).
…rt halves Host wiring (ds4.c): widen the Qwen GPU forward/session/dispatch guards from __APPLE__ to DS4_QWEN_GPU_BUILD (= __APPLE__ || ROCm build). There is no DS4_BACKEND_ROCM enum — the ROCm build runs under DS4_BACKEND_CUDA — so a new ds4_backend_is_qwen_gpu() helper replaces the backend==METAL checks in is_qwen35_metal / session-create / options_valid / resident-config. Engine open gates on DS4_QWEN_EXPERIMENTAL_ROCM=1 and forces resident (Qwen SSD streaming is not ported; the 19.4 GiB payload fits on the 64 GB box). The 1353-line forward block is backend-neutral (calls only ds4_gpu_*) and compiles unchanged. MoE (ds4_rocm_moe_launch.cuh): the shared ROCm Q4_K MoE core is specialized for DeepSeek's top-6 (fused sum6 kernels), but Qwen routes top-8. Run 8 as two top-4 halves + add, without touching the sum6 path: routed_moe_batch_tensor splits [n_token][8] into two [n_token][4] (new moe_split_selected_kernel) and adds the partials; the decode entry point uses the offset-0/16 half views. DeepSeek (n_expert=6) is unaffected. Measured on the Radeon 780M, resident: Qwen3.6-35B-A3B runs end-to-end on the GPU, coherent output, prefill 11.75 / generation 14.93 tok/s (vs ~10-11 CPU).
…1 parity End-to-end on the Radeon 780M (resident, 19.37 GiB): generation 13.4-14.9 tok/s (vs ~10-11 CPU reference), coherent output. CPU-vs-GPU parity: top-1 agreement 5/5, cosine 0.9994-0.9999 on top-k logits (evidence JSON attached). All 22 per-op kernel checks pass. Details, method, and next steps in RESULT_QWEN_ROCM.md.
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
1 participant
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.
Ciao Andrea! Seguito della PR #3. Il tuo
PREREG_QWEN36_35B_A3Bdice che laparity CUDA/ROCm servirebbe "~15-25 giorni più accesso all'hardware".
L'hardware ce l'ho — Qwen3.6-35B-A3B ora gira sul mio Radeon 780M (gfx1103,
ROCm 6.3.1) via il backend ROCm, resident, con output coerente e parità
numerica col tuo riferimento scalare CPU.
Risultato misurato (780M, resident, 19.37 GiB in GTT)
Batte già il tuo riferimento CPU scalare. Non batte ancora llama.cpp — atteso
per un v1 correctness-first (tutto F32, niente HIP-graph, MoE top-8 come due
top-4). I tuoi numeri Metal (58–66 t/s su M5 Pro) sono su ~2× la banda di
memoria e un set di kernel maturo.
Correttezza (gate M2, CPU-vs-GPU top-logprobs)
0.99994)
near-tie (comportamento upstream documentato)
tests/test_qwen35_rocm.c) passano vsds4_qwen_ref.c; il ricorrente Gated-DeltaNet combacia a ~1e-8Cosa contiene
rocm/ds4_rocm_qwen.cuh— 19 kernel HIP, port 1:1 dimetal/qwen35.metal(simdgroup Metal → wave32 RDNA3,
simd_sum/max/min→ butterfly__shfl_xor;gli stride in byte degli args ricalcano
ds4_metal.m) + i 20 entry pointds4_gpu_qwen35_*.specializzato per il top-6 di DeepSeek (kernel
sum6fusi). Qwen instradatop-8 → lo eseguo come due metà top-4 + add (
moe_split_selected_kernelrouted_moe_batch_two_halfinds4_rocm_moe_launch.cuh; il decode usa leviste half a offset 0/16). Il path DeepSeek (n_expert=6) è invariato.
ds4.c). Guardie del forward/session/dispatch Qwen allargateda
__APPLE__aDS4_QWEN_GPU_BUILD. Nota: non esiste enumDS4_BACKEND_ROCM— il build ROCm gira come
DS4_BACKEND_CUDA— quindi un helperds4_backend_is_qwen_gpu()sostituisce i checkbackend==METALinis_qwen35_metal / session-create / options_valid / resident-config. Il blocco
forward da 1.353 righe è backend-neutro e compila invariato. Gate env
DS4_QWEN_EXPERIMENTAL_ROCM=1; resident forzato (lo streaming SSD per Qwennon è portato — il payload sta in RAM sui 64 GB).
tests/qwen/normalize_qwen36_gguf.py— riproduce l'artefatto benedettodal sorgente Unsloth UD-Q4_K_S (Q6_K ffn_down_exps → Q4_K, output → Q8_0,
pad-id + template). Inventario risultante f32=361 / q8_0=252 / q4_k=120 ==
il tuo istogramma del prereg.
qui perché senza non compila).
Stato di verifica onesto
Compila e gira sul 780M; i numeri e la parità vengono da questo hardware. Il
diagnostico
--metal-graph-prompt-testè DeepSeek-only, quindi la parità è via--dump-logprobsCPU vs GPU (JSON intests/qwen/qwen36_rocm_cpu_parity_*.json).Non ancora fatto: soak 8k/32k +
ds4-benchCSV, throughput (HIP-graph, fusedgate/up+swiglu nel MoE a due metà, GQA-decode parallelo — ora è wired solo il
kernel GQA seriale), smoke del server su ROCm. Il server Qwen resta text-only
(niente tool-call) come da tuo runtime.
Dettagli completi in
RESULT_QWEN_ROCM.md. Marco ti scrive anche a parte —è più bello di una notifica GitHub :)
— Marco (con Claude), branch
marmyx77:rocm-qwen🤖 Generated with Claude Code