rocm: AMD Instinct (CDNA/gfx90a) support with multi-GPU device placement - #602
Open
ewindisch wants to merge 26 commits into
Open
rocm: AMD Instinct (CDNA/gfx90a) support with multi-GPU device placement#602ewindisch wants to merge 26 commits into
ewindisch wants to merge 26 commits into
Conversation
Adds `make mi200` / `make mi300` build aliases and DS4_ROCM_HAS_WMMA_W32 / DS4_ROCM_HAS_MFMA_W64 capability macros so device code can tell RDNA (wave32, WMMA) apart from CDNA (wave64, MFMA) at compile time, since host code has no target-arch macros to test. Also fixes the ROCm link itself on this toolchain: hipcc links PIE by default while this gcc compiles non-PIE by default, which fails at link time with R_X86_64_32 relocation errors unless the C objects are also built -fPIE. And since not every distro puts the ROCm lib directory on the default loader path, an rpath is added so the binaries don't require LD_LIBRARY_PATH. `make test-rocm-matmul-q8` is wired in here; its source lands in the next commit in this stack.
The batched Q8_0 GEMM used for large prefill chunks was built directly on __builtin_amdgcn_wmma_f32_16x16x16_f16_w32, an RDNA3/RDNA4 wave32 intrinsic that CDNA parts don't have (they use MFMA instead, and run wave64). This is why the plain `make rocm ROCM_ARCH=gfx90a` build failed outright with "needs target feature wmma-256b-insts,wavefrontsize32". Adds matmul_q8_0_f32_batch_mfma_4w_kernel: same 64x64 output tile and K_TILE=32 activation staging as the WMMA kernel, rebuilt on v_mfma_f32_16x16x16f16 with 4 waves of 64 threads instead of 4 of 32, each lane now dequantizing a quarter of the K range instead of all of it. The MFMA fragment layout (A/B/D lane mapping) was verified against a CPU reference on real gfx90a hardware before writing the kernel, not taken from documentation. Host dispatch (ds4_gpu_matmul_q8_0_tensor) picks between the two at runtime via cuda_device_wave_size(), since host-side code has no target-arch macro to test at compile time -- only the device passes do. An unrecognized wave size falls through to the existing portable sharedx kernel rather than launching a matrix-core kernel that would have compiled to a no-op for that target. tests/test_rocm_matmul_q8.c checks the new kernel against both the portable kernel and a CPU reference that models the f16 rounding the matrix core applies to its operands, run via `make test-rocm-matmul-q8 ROCM_ARCH=gfx90a`. On an MI210 this kernel matches the CPU f16 reference to 0.000% of output rms, and the portable kernel to 0.137% (the portable path stays in f32 activations, so that residual is expected f16 quantization, not a bug). Confirmed the test actually catches a layout error, not just quiet correctness, by deliberately perturbing one epilogue index and rerunning it: 491% error, immediately caught. RDNA (gfx1151) is unaffected; `make rocm ROCM_ARCH=gfx1151` still builds and takes the unchanged WMMA path.
The backend built its __shfl_*_sync masks from threadIdx.x, e.g. `0xffu << (threadIdx.x & 24u)` for an 8-lane group. On wave32 that happens to line up with the lane's position, but on wave64 threads 32..39 compute the same mask as lanes 0..7: the mask names the wrong half of the wave. HIP validates a shuffle's mask against the set of lanes actually reaching it, so this reliably aborted the queue with "HSA_STATUS_ERROR_EXCEPTION: An HSAIL operation resulted in a hardware exception" as soon as a kernel using one of these reductions ran (first hit: moe_gate_up_mid_decode_lut_qwarp32_kernel during decode, confirmed via rocgdb). Separately, five call sites did `__shfl_sync(FULL_WARP_MASK, x, 0)` with no explicit width, which defaults to `warpSize` -- 64 on CDNA -- so those broadcasts silently pulled from the wrong lane instead of lane 0 of the intended 32-wide group. Adds ds4_lane_in_wave() (wraps __lane_id() so it's correct regardless of hardware wave size) and ds4_subwave_mask<WIDTH>(), which builds the mask from the true lane position. DS4_WARP32_MASK replaces the old FULL_WARP_MASK constant at every one of these sites, and the five bare broadcasts get an explicit width of 32. On wave32 and on CUDA ds4_subwave_mask<32u>() reduces to the same all-ones constant as before, so this is a no-op there; the fix only changes wave64 behavior. Verified with rocgdb that the prior state crashed on `hi -n 4` mid-decode, and that the same run now completes cleanly through generation on an MI210 after this change.
…t kernels indexer_scores_wmma128_kernel assigns one 16x16 output tile per hardware wave (8 waves of 32 threads, one tile each). On a wave64 part the same 256-thread block only has 4 waves, so half the c_sh output tiles were silently left unwritten -- wrong scores, not a crash, since nothing about the wmma calls themselves is wave-size-specific once the tile loop covers the right range. Ports it to cover TILES_PER_WAVE = 8 / N_WAVES tiles per wave, computed from DS4_ROCM_HAS_MFMA_W64 (4 waves of 64, 2 tiles each on CDNA; unchanged 8 waves of 32, 1 tile each on RDNA). Separately, the hot-expert WMMA kernels in ds4_rocm_moe.cuh (used when a small number of experts see most of the batch) launch 32*MT-thread blocks and hand one matrix fragment to each 32-thread group -- a wave32-specific shape that would need reworking for CDNA's 64-wide waves, not just a launch-config tweak. Rather than port those now, gates use_wmma_hot and use_iq2_gate_wmma off when cuda_device_wave_size() != 32, so CDNA takes the portable expert-batch kernels already used in quality mode. Correct, if not yet the fastest possible path for the hot-expert case on CDNA.
…ing prefill
The SSD-streaming batched graph-prefill path (metal_graph_prefill_layer_major's
split_commands branch, the "gpu prefill layer N/M" progress messages) produces
output that is almost entirely independent of the actual prompt on wave64
hardware. Confirmed on an MI210 with --dump-logits and with real generation:
every broken run -- 2, 9, 10, 11, 15, 23, and 151 real tokens, several totally
different prompts -- converges on nearly the same top-5 vocab logits
(~token ids 1453/2435/3934/16291/57168), which is the signature of a
computation that stopped depending on its input, not a subtle numerical
divergence. Bisected the exact trigger against
metal_graph_use_streaming_decode_prefill's existing token-count cutoff: below
it, prefill runs one token at a time through the per-token decode-style path
(the "gpu streaming prefill token N/M" messages) and matches the CPU backend;
above it, prefill switches to the batched path and breaks, with the crossover
landing exactly on that existing cutoff (confirmed at both its default value
and by overriding it).
What this rules out, checked directly rather than assumed:
- Not the wave64 shuffle-mask fixes from the previous commit in this stack --
those made the crash go away; this is a different, silent bug that
predates them and survives them.
- Not the new MFMA matmul kernel -- gated to n_tok >= 256, never reached by
any of these reproductions.
- Not attention kernel choice -- forcing the batched path off the fast
attention_static_mixed_heads8_online_kernel and onto the hipBLAS GEMM
fallback (by disabling its dispatch condition and rebuilding) produces the
same broken output to three decimal places.
- Not g_quality_mode -- quality mode disables most of the opt-in fast
kernels and still produces the same broken logits on the same prompt.
- Not a wave-size assumption newly found in the "batch-only" kernel set
(compressor_*, hc_expand4*, head_rms_norm, the MoE tile-building kernels,
dequant/f32_to_f16) -- audited each for hardcoded 32-lane assumptions the
way the indexer WMMA kernel had; found none.
That leaves the ROCm-specific SSD-streaming expert-paging machinery unique to
that branch (rocm_graph_stream_layer_expert_load and the
layer_pagein/layer_readahead/layer_pread/batch_selected_addr machinery around
ds4.c:33096 onward) as the untested remainder. Isolating further needs
either a model that fits in one GPU's VRAM without SSD streaming (to test
whether the plain non-streaming batch path is also broken, which would move
suspicion back to something arch-generic) or per-layer intermediate-state
dumps compared against the CPU backend -- both out of reach on this hardware
in the time available.
The per-token streaming decode-prefill path is proven correct at every length
tested, including a 151-token prompt, so rather than ship a known-wrong batch
path, this makes the ROCm backend prefer streaming unconditionally on wave64
hardware: metal_graph_streaming_decode_prefill_max_tokens() returns
UINT32_MAX when ds4_gpu_wave_size() == 64, after the existing env-var override
(which still works, and still lets DS4_ROCM_DISABLE_STREAMING_DECODE_PREFILL=1
force the broken batch path back on for whoever picks up the real fix).
metal_graph_use_streaming_decode_prefill's separate quality-mode exclusion is
also lifted specifically on wave64, since quality mode goes through the same
broken batch path and produces the same garbage. Adds ds4_gpu_wave_size() to
the shared GPU API (ds4_gpu.h) with a real implementation in
ds4_rocm_runtime.cuh and a Metal stub returning 0, since ds4.c is compiled
once per backend and needs the symbol regardless of which backend provides it.
This costs prefill throughput on long prompts on CDNA -- streaming is
sequential per token where the batch path would parallelize -- but wrong
output is a worse failure mode than slow output. RDNA (Strix Halo, the
original ROCm target) is untouched: ds4_gpu_wave_size() returns 32 there, so
it keeps the batch path's original 18/64-token cutoff exactly as before.
Known gaps this does not cover, both unverified either way rather than
confirmed broken:
- Quality mode combined with SSD streaming hits an unrelated pre-existing
error during decode ("ROCm SSD streaming routed MoE missing compact
selected experts") that reproduces identically on the pre-this-commit
tree, so it is not new here, but it does mean quality mode + SSD
streaming does not fully work on this hardware regardless of this fix.
- The plain GPU-resident batch path without SSD streaming was never
exercised -- the only available model (80.76 GiB) does not fit in one
MI210's 64 GiB without streaming -- so whether the underlying bug is
SSD-streaming-specific or more general is genuinely unknown.
Validated end to end: real generation (not just --dump-logits) at 2, 10, 21,
and 151 real prompt tokens all produce coherent, on-task output matching the
CPU backend's register; make test-rocm-matmul-q8 still passes; make rocm
ROCM_ARCH=gfx1151 and make cpu both still build clean.
The ROCm multi-row fallback in ds4_gpu_hc_split_weighted_sum_norm_tensor collapsed every HC row, then passed the full output to the single-row rms_norm_weight_tensor entry point. Only row zero was normalized; prompt rows 1..N-1 remained zero in norm_out. Batched graph prefill therefore stopped depending correctly on most prompt tokens, producing attenuated, largely prompt-independent logits in both multi-GPU mode and the SSD batch path. Use ds4_gpu_rms_norm_weight_rows_tensor with the row count derived from the output tensor, matching the correction already present in the CUDA implementation. The layer-0 last-prompt-row attn_norm changed from all-zero to bit-identical with the streaming reference; greedy output changed from 特 to We. Verification on two MI210s (gfx90a): make mi200 and test-rocm-matmul-q8 pass. Twenty independent multi-GPU runs all exited 0, produced all 43 requested dumps, generated the same 'We need' continuation, and were byte-identical at every layer. Worst norm ratio versus the single-GPU reference was 0.277 at layer 17; layer 40 was 0.258, versus the prior 340-1500x failures.
metal_graph_debug_dump_tensor's dump hook always read row 0 of a [n_tokens][per_row] batch tensor. Batched graph prefill evaluates all prompt rows in one call at position 0, so comparing a specific token's state (e.g. the final prompt row) required this row selector -- without it, only row 0 was ever inspectable, masking exactly the kind of last-row-specific corruption this PR's root cause turned out to be. Adds DS4_ROCM_GRAPH_DUMP_ROW (byte offset = row * n_f32 * 4), env-gated, no effect when unset.
The CDNA SSD-streaming batch-prefill workaround was masking the same batched HC normalization bug fixed in the preceding commit. With every row now normalized, forced batch-prefill logits are prompt-dependent and correct-shaped again for 14-, 16-, and 99-token prompts, so remove the wave64 always-streaming override and restore the original token cutoff and quality selection. This also fixes a separate hang exposed after correct batch prefill. The final layer may leave an asynchronous selected-expert read in the shared read-job pool because no next layer exists to consume it. Hotlist seeding immediately reuses that pool and previously waited forever on the stale 147-job batch. Drain both batch and single selected-read state when releasing the temporary layer cache. Validated on MI210/gfx90a with the 80.76 GiB Flash model: three forced batch-prefill dump-logits cases (14, 16, and 99 prompt tokens) all exit successfully; the first-token argmaxes are prompt-dependent (We, I, We), rather than the old content-independent token 1453 signature. After restoring the default cutoff, two >64-token layer-major SSD-prefill runs complete real greedy generation with coherent prompt-specific continuations at 2.17-2.29 prefill tok/s and 1.26-1.28 generation tok/s. make mi200 succeeds. Known gap retained honestly: --quality with SSD streaming still reaches the pre-existing decode-time 'missing compact selected experts' error after successful batch prefill. This commit does not claim to repair that independent quality-mode cache-selection issue.
Ports ROCm's single-device-only limitation away, mirroring ds4_cuda.cu's
existing "GPU-only multi-tier" pipeline-placement scaffolding: per-device
streams/hipBLAS/hipBLASLt handles, a per-device selective weight cache
wired into cuda_model_range_ptr (the single chokepoint every existing
kernel launcher already funnels weight lookups through, so none of the
~200 launchers needed per-call edits), and cross-device tensor copy with
peer-access validation plus a pinned-host-bounce fallback (peer access
between the two MI210s on this box is unavailable -- PCIe-only, no
bridge/XGMI). Does not implement --cuda-tensor-parallel (row/col-sharded
compute split live across two tiers within one layer); that stays
CUDA-only exactly as before. New file rocm/ds4_rocm_mgpu.cuh; struct
ds4_gpu_tensor gained a device_id field (see that file's header comment
for why the layout has to stay byte-identical across ds4_rocm.cu and
ds4_rocm_compat.cu, which reach it via different routes).
Six real bugs found and fixed this session, each confirmed against a live
dual-GPU run of an 80GB model split across both cards (--gpu-devices 0,1
--gpu-vram auto, no --ssd-streaming; GPU0: layers 0-28+embedding, GPU1:
layers 29-42+output head):
1. System prereq, not a code bug: this box's ulimit -l (locked memory)
was 8 MiB, far too small for ROCm's DMA staging buffers across two
devices -- caused a driver-level pin-buffer exhaustion flood
("DmaBlitManager::getBuffer failed to pin a resource") followed by a
hard crash in libamdhip64 on the next unrelated kernel launch. Fixed
by raising memlock to unlimited via /etc/security/limits.d/ (requires
a fresh login session; PAM limits apply at login, not retroactively).
2. ds4_gpu_register_model_map_no_copy was hipHostRegister-ing the whole
~80 GiB model with hipHostRegisterMapped (SVM/zero-copy). With two GPU
agents registered this overran the driver's resident-SVM working-set
limit ("amdgpu: SVM mapping failed, exceeds resident system memory
limit" flood, then a crash). Removed entirely -- not needed, since the
per-device selective cache uploads every GPU-placed tensor before any
layer runs, so the "uncached fallback" zero-copy pointer this
registration existed for is never on the hot path here.
3. Performance, not correctness: after fix #2, weight loading fell back
to demand-paged cudaMemcpy straight off the mmap -- ~15 MiB/s despite
this box's NVMe RAID1 doing ~2 GiB/s on a plain sequential dd (and
PCIe genuinely running x16 Gen4, confirmed independently via
rocm-validation-suite's PEBB module at ~25.8 GB/s). Fixed with a
synchronous bulk sequential pread() to warm the page cache before the
per-tensor copy loop, plus a single reusable pinned staging buffer for
the H2D copies instead of letting HIP re-pin a bounce buffer on every
one of ~1300 calls per tier. Full 80 GiB load: 20+ minutes -> well
under a minute.
4. g_cuda_tmp (cuda_tmp_alloc's scratch buffer, used e.g. for F32->F16
activation conversion before a cuBLAS/hipBLASLt GEMM) was a single
global buffer shared across both tiers -- a tier-1 kernel could get
handed a tier-0-resident pointer and fault ("Memory Fault Error ...
kernel: f32_to_f16_kernel"). Fixed by making it per-tier.
5. cuda_q8_f16_ptr / cuda_q8_f16_transpose_ptr (the Q8->F16 weight-dequant
cache backing the cuBLAS/hipBLASLt fast matmul path) had the same bug
class: cached purely by (model_map, offset, weight_bytes, in_dim,
out_dim) with no device dimension, so it could hand back another
tier's device-resident dequant buffer -- silently wrong matmul output
rather than a fault, since the returned pointer is a valid-looking
VRAM address just on the wrong device. Added a device_id field to both
cache-entry structs and filter every lookup on the active device.
6. Added a same-device guard before launching f32_to_f16_kernel in
ds4_gpu_tensor_copy_f32_to_f16, matching CUDA's own reference
(ds4_cuda.cu already refuses rather than launching cross-device here).
Net result: the dual-GPU run now completes end-to-end with no crash
(prefill + generation both finish, ~2-3 tok/s). Generation output is
still wrong, though: the first token is garbled, later decode tokens read
as locally-plausible English continuations -- traced live with a new
DS4_MGPU_TRACE=1 env var in metal_graph_set_active_tier_decode/_batch
(ds4.c), no failed cross-tier copies, tier-switch pattern looks correct,
pointing away from the decode-to-decode hidden-state handoff and toward
the prefill-stage output-head/final-logits computation instead. Also
notable: --ssd-streaming refuses to combine with multi-GPU placement, so
this run exercises the plain (non-streaming) batch graph-prefill path,
which single-GPU builds never reached for this 80GB model (always forced
onto the streaming decode-prefill workaround from
ewindisch/mi210-05-prefill-correctness to fit in one card's VRAM). Still
open whether this is the same defect PR5 worked around or a new one in
code ROCm had simply never exercised before. See
project_ds4_rocm_cdna_port.md memory for full debugging notes and next
steps (a smaller model that fits one MI210 without streaming would let
the two be compared directly).
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Same bug class as the g_cuda_tmp and q8_f16-dequant-cache fixes in the
previous commit: g_shared_gate_up_stream/_ready_event/_tmp/_pending were
single global instances reused across tiers. The stream and event are
device-bound (created while one tier's device was active, invalid to
launch/sync against once a different tier's device is current), and the
scratch buffer + "pending" flag didn't distinguish "tier 0 has an
outstanding async op" from "tier 1 does" -- a tier-1 call could reuse
tier 0's stream (invalid cross-device stream use) and/or race a
wait/reuse against the wrong tier's in-flight op. Indexed all five by
g_current_logical_tier, matching g_cuda_tmp.
Confirmed via a live dual-GPU prefill run (--dump-logits) that this
specific function isn't exercised during prefill (bit-identical argmax
before/after: id 1453 "特", logit 12.522047) -- it's evidently
decode-only, and prefill's shared-expert gate/up goes through the
already-tier-safe synchronous ds4_gpu_shared_gate_up_swiglu_q8_0_batch_tensor
instead (default stream, no cross-call caching). Keeping the fix anyway:
it's a real latent bug for decode once the placement work in the
previous commit's multi-tier prefill correctness issue is resolved, and
this file is otherwise dead code with respect to that investigation.
The multi-tier batch-prefill correctness bug documented in the previous
commit remains open: exhaustively audited every static/global cache in
the rocm/*.cuh kernel launchers (this was the last of three: g_cuda_tmp,
cuda_q8_f16_ptr/_transpose_ptr, and this file -- no others exist outside
the SSD-streaming-specific machinery, which isn't active in this
scenario). Also directly ruled out the "portable q8 kernels" matmul
fallback path being buggy in isolation: forcing single-GPU onto that
exact code path via DS4_CUDA_NO_Q8_F16_CACHE=1 still produces the correct
top token ("We", logit 34.13 vs the cached path's 36.64 -- an expected/
normal precision difference, not a bug). The prefill logits correlation
between single-GPU and multi-GPU stays weak (~0.21) with a compressed
dynamic range on multi-GPU (max logit 12.5 vs 36.6), suggesting an
attenuated/partially-missing contribution rather than pure corruption --
but the source of that attenuation is still unlocated. See
project_ds4_rocm_cdna_port.md for full notes; next productive step is
likely per-layer intermediate-state comparison or a smaller model that
fits one MI210 without SSD streaming, to get a single-GPU baseline on
the exact same non-streaming batch-prefill code path.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Same bug class as the previous two commits' g_cuda_tmp and
g_shared_gate_up_* fixes: g_q8_f16_bytes and g_q8_f16_disabled_after_oom
were single global instances shared across tiers.
cuda_q8_f16_cache_release_all() freed every cached range's device_ptr
without cudaSetDevice()'ing to that range's owning device first --
cudaFree on a pointer belonging to a different device than the current
one doesn't actually free it, leaking VRAM and desyncing the accounting.
And because the disable-after-OOM flag and byte counter were global, one
tier hitting its (much tighter, since multi-GPU weights already consume
most of that tier's budget) cache ceiling would disable and release the
OTHER tier's still-valid, still-in-use cache too. Added
cuda_q8_f16_cache_release_tier(tier) so an OOM on one tier only tears
down that tier's own entries, and indexed the byte counter and disable
flag by tier via a new cuda_q8_f16_tier() helper (mirrors
cuda_shared_gate_up_tier() from the previous commit).
Also added DS4_Q8F16_TRACE (env-gated, matches the existing
DS4_MGPU_TRACE convention): logs every q8_f16_ptr/_transpose_ptr cache
store (label, offset, dims, active device, tier) and every
cuda_model_range_ptr strict-lookup miss that falls through to the legacy
non-device-aware range cache.
This does NOT fix the multi-tier batch-prefill correctness bug from the
previous two commits -- confirmed via a live dual-GPU per-layer
hc_ffn_post dump comparison against the known-good single-GPU baseline.
Investigated and ruled out as the cause of that bug specifically:
- Cross-device q8_f16 cache hits: the per-device identity check
(device_id == active_device) never missed in a traced run (0 STRICT
MISS events), and the lazy caching that happens during real
inference (ds4_rocm_matmul.cuh etc. call cuda_q8_f16_ptr directly,
not just the ds4.c preload sweep) correctly picks up each tier's own
device context via the existing ds4_gpu_set_current_device plumbing.
- Dequant content corruption: dumped the first cached tensor's
dequantized f16 values and the raw q8_0 source bytes and compared
against a manual CPU-side Q8_0 decode -- bit-exact match.
- The q8_f16 cache as sole cause, full stop: forcing
cuda_q8_f16_cache_allowed() to always return 0 under g_n_gpus > 1
(i.e. fully disabling this cache, not just letting budget deny it)
still reproduced the same layer 5->6 divergence blowup in the
hc_ffn_post comparison, so this commit does not attempt that -- it
would cost the acceleration for no correctness benefit.
- A real but separate bug found along the way: the one-shot
"attention output warmup" GEMM (cuda_q8_f16_warmup_attention_output_
a/b_gemm) and the preload sweep's optional_q8_preload_disabled latch
in ds4_gpu_cache_q8_f16_range are both plain (non-per-tier) static
locals -- the latter means the FIRST tensor the preload sweep is
denied on (tier 0, budget-limited almost immediately under
multi-GPU) permanently stops the sweep from even trying the rest of
tier 0 or any of tier 1, though tier 1 still gets cached lazily
during real inference via the matmul-launcher call sites. Neither
explains the bug on its own (see above).
New finding worth recording for the next session: the divergence is
NOT deterministic. Two back-to-back runs of the identical binary,
config, and prompt produced different generated tokens and different
per-layer divergence magnitudes/onset (layer 5 norm-ratio 2.4x then
8.9x vs. baseline across two runs; layer 0-4 stayed stable and
reproducible both times). That points at a race or scheduling-order
sensitivity somewhere in the multi-tier dispatch rather than a static
data/logic bug, which the extensive static-cache auditing across this
and the previous two commits was aimed at. Worth auditing next:
function-local "run once" static latches other than the two named above
(grep for `static int ... = 0;` inside rocm/*.cuh), and whether any
async op related to first-use-per-layer state (MoE routing/hotlist,
indexer, attn_compressor -- all introduced or reintroduced within
layers 2-6 of this model's architecture, unlike layers 0-1) is missing a
synchronization point that only sometimes loses the race.
See project_ds4_rocm_cdna_port.md for full session notes.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Accumulated debugging instrumentation from the multi-GPU correctness investigation, kept for future ROCm debugging since it's genuinely useful and costs nothing by default: - DS4_FFN_FORCE_SYNC: forces a full device sync immediately after routed-MoE and after shared-expert computation, every layer. Used to test for a same-GPU kernel-ordering race at the routed-MoE/shared-expert boundary (result: negative -- ruled that out as a contributing factor). - DS4_FORCE_SYNC_LAYER / ds4_moe_layer_force_sync(): scopes a full device sync to a specific layer index, used to bracket every kernel in routed-MoE's live dispatch path at layer 40 specifically (result: negative -- ruled out a kernel-ordering race within routed-MoE at that layer). - ds4_moe_integrity_hash64 / ds4_moe_integrity_dump_device / ds4_moe_integrity_compare_selected_weights: dumps and FNV-1a-hashes layer-scoped intermediate tensors and cross-checks device-resident selected-expert weight ranges against the model mmap, to distinguish a kernel receiving a discrete stale/wrong input or weight image from a genuine compute bug. Used during the final integrity-verification pass that helped confirm the row-normalization root cause. All gated behind env vars, no effect when unset. None of this is the fix itself -- see the row-normalization and SSD-prefill-drain fixes for that.
make mi200/make mi300 targets, DS4_ROCM_HAS_WMMA_W32/DS4_ROCM_HAS_MFMA_W64 arch macros, -fPIE + rpath link fixes for this gcc/hipcc combo.
MFMA counterpart of the RDNA-only WMMA batched Q8_0 matmul kernel -- the thing that made gfx90a fail to compile at all originally. Validated by tests/test_rocm_matmul_q8.c against a CPU f16 reference.
Fixes __shfl_*_sync masks built from threadIdx.x (wrong on wave64, right by coincidence on wave32) and 5 sites with an implicit warpSize-width broadcast. Fixed an actual HSA_STATUS_ERROR_EXCEPTION hardware-exception crash.
…otlist kernels Ports indexer_scores_wmma128_kernel's per-wave tile assignment to wave64 (was silently leaving half the output tiles unwritten on CDNA). Gates two wave32-shaped 'hotlist' MoE WMMA kernels off on wave64, falling back to the portable kernels quality mode already used.
Root cause: ROCm's batched HC fallback (ds4_gpu_hc_split_weighted_sum_norm_tensor) collapsed every prompt row's HC contribution correctly, but then normalized only row 0 via the single-row rms_norm_weight_tensor entry point, leaving norm_out zero for every other row -- so batched graph prefill produced attenuated, largely prompt-independent logits on wave64, in both the SSD-streaming batch path and (as later confirmed) native multi-GPU placement. Fix: use ds4_gpu_rms_norm_weight_rows_tensor with the row count derived from the output tensor, matching the correction already present in the CUDA implementation. Also restores full batched SSD-streaming prefill throughput by removing the prior wave64 always-streaming workaround now that the actual defect is fixed, and drains a final-prefill async selected-expert read state that could otherwise hang after correct batch prefill. Verified: 20 independent multi-GPU runs byte-identical at every layer; forced batch-prefill dump-logits correct and prompt-dependent at 14/16/99 tokens; two real >64-token generations complete coherently at ~2.2 prefill tok/s (vs. ~1.0 under the old workaround) and ~1.27 generation tok/s. Known separate gap, not claimed fixed here: --quality --ssd-streaming still hits a pre-existing 'missing compact selected experts' decode-time error.
Splits pipeline/layer placement across GPU tiers for ROCm (mirroring ds4_cuda.cu's existing CUDA-only multi-tier scaffolding), letting an 80GB model that doesn't fit on one MI210 run split across two without needing --ssd-streaming. New rocm/ds4_rocm_mgpu.cuh: per-device streams/hipBLAS/ hipBLASLt handles, a per-device selective weight cache wired into the one chokepoint every kernel launcher already used, and cross-device tensor copy with peer-access validation + pinned-host-bounce fallback (no P2P bridge between these two MI210s). Three real per-tier-globals bugs found and fixed along the way -- global scratch/cache/stream state shared across tiers instead of indexed per device, silently handing one tier's device-resident pointer to the other: the F32->F16 conversion scratch buffer, the shared-expert async gate/up state, and the Q8->F16 dequant cache's disable/free accounting. The actual multi-GPU correctness bug (nondeterministic-looking garbled output, ~40-60% of runs) turned out to share PR5's root cause exactly -- see that merge for the fix. An extensive but ultimately unnecessary mitigation (gating 9 cuBLAS GEMM call sites off under multi-GPU due to suspected runtime algorithm nondeterminism) was tried, found to not measurably affect the failure rate either way once properly A/B tested, and has been dropped from this branch's history now that the real cause is fixed, to avoid paying its performance cost for no benefit. Verified: 20/20 repeated multi-GPU runs deterministic and byte-identical at every layer, correct 'We need' continuation matching the single-GPU reference.
optional_q8_preload_disabled was a single static int shared across tiers, so a tier-0 budget denial (tier 0 is VRAM-tight under multi-GPU) permanently stopped the preload sweep from ever trying tier 1, which may still have headroom. Made it per-tier (g_q8_f16_preload_disabled[ DS4_MAX_GPUS]), matching the existing per-tier convention used by g_q8_f16_disabled_after_oom.
Both matmul_q8_0_pair_f32_sharedx_warp_rows_w32_kernel and moe_gate_up_mid_decode_lut_qwarp32_kernel were tiled for 32-lane wavefronts (inherited from this project's original RDNA3/Strix-Halo- only codebase), made correct on wave64/CDNA via PR3's sub-wave shuffle-mask partitioning but never reshaped for wave64 efficiency -- each 64-lane wavefront ran two time-multiplexed 32-wide passes instead of one native 64-wide one. Kernel-level profiling on a 100-token decode run showed these two kernels alone accounted for 32.1% of all GPU kernel execution time. Adds a wave64-native counterpart of each, dispatched via the existing cuda_device_wave_size()==64 gating idiom (same pattern PR2 used for the CDNA MFMA batched Q8_0 kernel). The original w32/qwarp32 kernels are completely untouched, so RDNA3/Strix Halo sees zero behavior change. matmul_q8_0_pair_f32_sharedx_warp_rows_w64_kernel: instead of two independent 32-lane sub-groups each computing one row (PR3's correctness fix), all 64 lanes cooperate on a single row -- lanes 0-31 consume even Q8_0 blocks, lanes 32-63 consume odd blocks in the same iteration, combined by one native 64-wide shuffle-reduction. Rows per block halved (32->16) with grid doubled to compensate. 46% reduction in average per-call kernel time (1166.9us -> 630.1us) on a 100-token decode profile. moe_gate_up_mid_decode_lut_hwarp32_kernel: widens the cooperating-lane group per row from 8 lanes (quarter_warp_sum_f32) to 16 lanes (half_warp_sum_f32) for this model's confirmed xq_blocks==16 decode shape, again halving rows-per-block and doubling the grid to preserve CU occupancy. 24.5% reduction in average per-call kernel time (1599.2us -> 1207.6us). Both validated against a CPU double-precision reference and cross-checked bit-exact/near-bit-exact against their untouched w32 counterparts (tests/test_rocm_matmul_pair_q8.c, tests/test_rocm_moe_gate_up_mid_hwarp32.c), using non-block-aligned dimensions to stress tail handling. Three other w32-shaped hotspot kernels from the same profiling pass (matmul_q8_0_hc_expand, matmul_q8_0_f32_sharedx_warp_rows, grouped_q8_0_a) were attempted with the same treatment and did not pan out -- two showed clear regressions on independent structural retries (root cause: the originals already achieve full wavefront occupancy via two independent parallel rows, or the win doesn't transfer to their extra per-row scalar epilogue), and one showed only noise-level improvement. None of those are included here.
…current before PR
moe_down_sum6_qwarp32_kernel used 8-lane cooperation (quarter_warp_sum_f32) per output row -- structurally identical in shape to moe_gate_up_mid_decode_lut_qwarp32_kernel before its own wave64 fix (also 8-lane originally), which won 24.5% after widening to 16-lane cooperation and halving rows-per-block / doubling the grid to preserve CU occupancy. This kernel has a trivial single-scalar-store epilogue and does not already use both wavefront halves for two independent full-width rows, so it matches neither regression pattern this session established. Adds moe_down_sum6_hwarp32_kernel applying the same fix: 16-lane cooperation (half_warp_sum_f32), rows-per-block halved 32->16 with the grid doubled to match. Dispatched via the existing cuda_device_wave_size()==64 gating idiom at both routed-MoE down-projection call sites, behind an opt-out (DS4_ROCM_DISABLE_MOE_DOWN_SUM6_HWARP32); the original qwarp32 kernel is completely untouched, so RDNA3/Strix Halo sees zero behavior change. Unlike the gate/up/mid fix this kernel has no "% 16 == 0" dispatch gate, so it must remain correct for any midq block count. 6.5% reduction in average per-call kernel time (579.1us -> 541.5us) on a 100-token decode profile, measured via rocprofv3. Validated against a CPU f64 reference and cross-checked near-bit-exact against the untouched qwarp32 kernel (tests/test_rocm_moe_down_sum6.c), using non-16/32-aligned dimensions (out_dim=777, midq_blocks=22) to stress the tail of the per-lane strided accumulation loop. Adds a test-only direct entry point (ds4_gpu_test_moe_down_sum6_tensor) to drive either kernel in isolation. Co-Authored-By: Claude <noreply@anthropic.com>
ds4_gpu_matmul_f16_tensor()'s ordered_decode path dispatches matmul_f16_ordered_chunks_kernel, which the w32 build launches as a 32-thread block. On a native-64-lane wavefront (CDNA/gfx9xx) that leaves the other 32 lanes of every dispatched wavefront completely idle -- not a two-independent-rows tradeoff like the other w32 kernels this session touched, just unused capacity. Adds matmul_f16_ordered_chunks_w64_kernel, which divides the same row's dot product across all 64 threads instead of 32, halving the per-lane chunk length. The reduction is still done by shared-memory store plus a single-thread sequential sum (matching the original's approach, not a shuffle reduction) to keep this a minimal, low-risk change isolated to the launch width. Dispatched via the existing cuda_device_wave_size()==64 gating idiom, behind an opt-out (DS4_ROCM_DISABLE_F16_ORDERED_W64); the original w32 kernel is completely untouched, so RDNA3/Strix Halo sees zero behavior change. 19.3% reduction in average per-call kernel time (346.9us -> 280.0us) on a 100-token decode profile, measured via rocprofv3. Validated against a CPU f64 reference and cross-checked near-bit-exact against the untouched w32 kernel (tests/test_rocm_matmul_f16_ordered.c), using the real router projection shape (in_dim=4096, out_dim=256). Adds the test-rocm-matmul-f16-ordered Makefile target to build/run it. Co-Authored-By: Claude <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
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.
Adds AMD Instinct (CDNA2, gfx90a — tested on dual MI210) support to the existing ROCm backend, which previously targeted RDNA3/Strix Halo only, plus native multi-GPU device placement and a set of wave64-specific performance fixes. Squashed history below is preserved as a series of well-scoped commits (each originally developed/tested as its own branch); happy to split into separate PRs if preferred, though GitHub doesn't support true cross-fork stacking so this landed as one PR — see commit-by-commit history for incremental review.
What's included
make mi200/make mi300targets,DS4_ROCM_HAS_WMMA_W32/DS4_ROCM_HAS_MFMA_W64arch macros, link fixes for this toolchain.make test-rocm-matmul-q8).__shfl_*_syncmasks built fromthreadIdx.xwere wrong on wave64 (right by coincidence on wave32); fixes the actualHSA_STATUS_ERROR_EXCEPTIONcrash on real hardware.matmul_q8_0_pair_f32_...,moe_gate_up_mid_decode_lut_...) were tiled for 32-lane wavefronts (inherited from the RDNA3-only codebase) and, while correct on wave64 via commit 3's fix, left roughly half of CDNA's per-wavefront throughput unused. Adds true 64-lane-native counterparts, dispatched via a runtime wave-size check, with the original kernels left completely untouched (zero behavior change on RDNA3/Strix Halo). 46% and 24.5% reductions in average per-call kernel time respectively, measured via rocprofv3 kernel-level profiling on a 100-token decode run. Both validated against a CPU double-precision reference with non-block-aligned dimensions to stress tail handling.Testing
make test-rocm-matmul-q8,make test-rocm-matmul-pair-q8,make test-rocm-moe-gate-up-mid-hwarp32all pass on gfx90a.--gpu-devices 0,1 --gpu-vram auto.g_n_gpus/device-count checks) — RDNA3/Strix Halo and single-GPU code paths are unmodified.Known gap, not addressed by this PR
Quality mode + SSD streaming hits a pre-existing "missing compact selected experts" decode error under ROCm. This reproduces independently of everything in this PR and is being tracked separately.