R9700 (Thunderbolt eGPU) as hip:0 source: hipMemcpyPeer SIGSEGV, blocks dual-GPU DS4 expert parallelism on Strix Halo pairing

[chrisjohnson]

Summary

hipMemcpyPeer segfaults (SIGSEGV, exit 139) specifically when an R9700 attached via a Thunderbolt/USB4 eGPU dock is the source device (hip:0) copying to an integrated Strix Halo APU (hip:1). The reverse direction (Strix Halo as source) works cleanly every time. This blocks the dual-GPU asymmetric expert parallelism DS4 path (dflash_server with DFLASH_DS4_MOE_TP_INPROC=1), which requires the discrete GPU as the initiating hip:0 device.

Isolated with a minimal ~90-line standalone HIP program — zero Lucebox code — so this looks like a box/driver-level HIP peer-access defect specific to this eGPU/iGPU pairing, not an application bug. Filing here first since the repro was found while qualifying this exact hardware pairing against DS4.md's documented setup, and to check whether your own reference R9700+Strix Halo machine uses a native PCIe slot or an eGPU dock — that topology difference may be the actual variable.

Environment

• AMD Ryzen AI Max+ 395 (Strix Halo), gfx1151, 128GB unified memory
• AMD Radeon AI PRO R9700, gfx1201, 32GB GDDR6, attached via a Thunderbolt/USB4 eGPU dock (not a native PCIe slot)
• ROCm 6.4.1, HIP backend, ghcr.io/luce-org/lucebox-hub:rocm image (also reproduced with a bare hipcc-compiled standalone program, no container-specific dependency)
• Dual-arch dflash_server build: -DDFLASH27B_HIP_ARCHITECTURES='gfx1151;gfx1201' -DGGML_HIP_GRAPHS=ON -DCMAKE_BUILD_TYPE=Release
• Models: DeepSeek-V4-Flash-ROCMFP2-STRIX.gguf (102.3GB) + DeepSeek-V4-Flash-DSpark-draft-Q4RMFP4-denseF16.gguf (11.3GB), matching DS4.md's documented "qualified R9700 + Strix Halo machine" setup exactly

What happens with dflash_server

With container HIP device ordering corrected so the R9700 lands on hip:0 (HIP_VISIBLE_DEVICES=1,0 — see note below on why this isn't 0,1), the server gets well into real placement logic:

[deepseek4] hybrid placement: gpu_total=31.86 GiB gpu_free=28.14 GiB core=3.72 GiB kv=0.89 GiB
  warm=0.25 GiB safety=0.50 GiB expert_budget=11.43 GiB hot/layer=32
[deepseek4] placement expert_memory: total=91.38 GiB uniform_round=365.50 MiB hot=1376 11.42 GiB
  cold=9632 79.95 GiB
[deepseek4-moe-tp] peer access hip:0 <-> hip:1: enabled

— then segfaults immediately after peer access is reported enabled, before any cross-device data transfer completes and before any measurable VRAM increase. Reproduced twice, identically: once via serve_ds4_dual_rocm_128k.sh, once invoking dflash_server directly with DS4.md's minimal-activation env vars (omitting --peer-access didn't help — DFLASH_DS4_MOE_TP_INPROC=1 triggers cross-device access internally regardless of that flag).

No stack trace on stdout (release build, no signal handler). coredumpctl shows SIGSEGV inaccessible for dflash_server at both crash timestamps. GPUs are fully healthy immediately after every crash (rocm-smi --showuse 0% busy both, no fault/IOMMU/DMAR entries in journalctl -k).

Minimal repro, no Lucebox code

To rule out an application bug, wrote a ~90-line standalone HIP C++ program: hipGetDeviceCount → per-device hipGetDeviceProperties → hipDeviceCanAccessPeer both directions → hipSetDevice/hipDeviceEnablePeerAccess both directions → hipMalloc on each device → hipMemcpyPeer (1MB) → exit, with fflush after every step. Compiled with hipcc --offload-arch=gfx1151 --offload-arch=gfx1201.

Each ordering run twice (4 runs total), 100% reproducible both times:

• HIP_VISIBLE_DEVICES=1,0 (R9700 on hip:0, Strix on hip:1): crashes, SIGSEGV, both runs. Device enumeration, hipDeviceCanAccessPeer (both directions report 1), hipSetDevice, and both hipDeviceEnablePeerAccess calls (both report success) all print cleanly. Crash happens inside/immediately around hipMemcpyPeer(d1 <- d0, 1MB) — the last line printed both times was Attempting hipMemcpyPeer(d1 <- d0, 1MB)...; the result line never appears.
• HIP_VISIBLE_DEVICES=0,1 (Strix on hip:0, R9700 on hip:1): succeeds cleanly, exit 0, both runs. hipMemcpyPeer -> no error and TEST COMPLETED WITHOUT CRASHING. both print.

Both orderings report hipDeviceCanAccessPeer = 1 in both directions and both report hipDeviceEnablePeerAccess succeeding in both directions — the crash isn't caught by either capability check, only by an actual hipMemcpyPeer data transfer, and only when the R9700 initiates it.

coredumpctl confirms the identical SIGSEGV inaccessible signature as the three dflash_server crashes above, at the exact timestamps of the two HIP_VISIBLE_DEVICES=1,0 runs.

Note: container HIP device ordering is inverted from host rocm-smi ordering (unrelated but adjacent gotcha)

Host-level rocm-smi --showproductname shows GPU[0] = R9700, GPU[1] = Strix Halo. Inside a container, KFD topology node numbering puts Strix Halo at HIP device 0 and R9700 at HIP device 1, regardless of --device bind-mount order — this is a kernel-global property. So HIP_VISIBLE_DEVICES=0,1 (matching host-level GPU numbering) actually puts the 128GB unified Strix Halo card on hip:0, not the R9700 the DS4 script expects there — worth calling out in DS4.md if it isn't already, since it produces an immediate, different, and more confusing crash (right after model metadata, before real placement logic) than the peer-DMA bug above.

Question for maintainers

Is your own qualified R9700 + Strix Halo reference machine using a native PCIe slot for the R9700, or an eGPU dock (Thunderbolt/USB4/OCuLink)? DS4.md doesn't say. My working hypothesis is that peer DMA between a Thunderbolt-tunneled discrete GPU and an integrated GPU is an architecturally unusual path that may not be exercised by your own qualification hardware if it's native PCIe — happy to help narrow this down further (e.g. testing with DFLASH_DS4_MOE_TP_INPROC disabled, or with a smaller allocation) if useful, or to move the R9700 to a native M.2 slot if one becomes available, to confirm whether topology is the actual variable.

Full writeup with additional log excerpts and methodology available on request.

[chrisjohnson]

Update: found the real crash site, and confirmed a fix — ROCm 7.2.2 resolves it entirely

Follow-up to the report above. Two findings since filing:

1. The crash isn't actually in hipMemcpyPeer/peer-DMA — a gdb backtrace on the real dflash_server crash (not just the minimal repro) shows it's in ggml_backend_cuda_buffer_set_tensor, a plain host→device weight load during build_moe_hybrid_storage_from_file (mmap'd GGUF → GPU), not a cross-device peer copy:

Thread 1 "dflash_server" received signal SIGSEGV, Segmentation fault.
0x000076aeb40e3b72 in ?? () from /opt/rocm/lib/libamdhip64.so.6
#0  0x000076aeb40e3b72 in ?? () from /opt/rocm/lib/libamdhip64.so.6
#1  0x000076aeb3f26ac1 in ?? () from /opt/rocm/lib/libamdhip64.so.6
#2  0x000076aeb3f28b53 in ?? () from /opt/rocm/lib/libamdhip64.so.6
#3  0x000076aeb3f39bde in ?? () from /opt/rocm/lib/libamdhip64.so.6
#4  0x000076aeb5d6581a in ggml_backend_cuda_buffer_set_tensor(...) from libggml-hip.so.0
#5  build_moe_hybrid_storage_from_file(...)
#6  build_moe_hybrid_storage_from_file_with_mmap(...)
#7  build_deepseek4_moe_hybrid_storage_from_file_with_mmap(...)
#8  DeepSeek4Backend::init_hybrid_model()

The kernel log shows a 100% deterministic crash signature — same fault address (segfault at 18, i.e. dereferencing near-NULL) and same code offset inside libamdhip64.so.6.4.60401 every single time, regardless of --peer-access, HSA_ENABLE_SDMA=0, or building ggml-cuda with -DGGML_CUDA_NO_PEER_COPY=ON (none of which changed anything — same crash, same offset). It also still crashes with --peer-access omitted entirely (so g_peer_access_opt_in never gets set and enable_peer_access_pair is never called) — the crash happens purely from the hybrid weight-loading path itself, independent of any application-level peer-access opt-in.

Given none of those knobs mattered, this looks like an internal HIP runtime bug tied to having both devices (R9700 + Strix Halo) initialized together in this specific topology, not something triggered by any particular copy pattern.

2. Confirmed fix: rebuilding against ROCm 7.2.2 (up from 6.4.1) resolves it completely. Same source checkout, same HIP_VISIBLE_DEVICES=1,0, same launch flags (including --peer-access) — installed rocm-dev=7.2.2.70202-86~22.04 + hipcub-dev7.2.2 from AMD's apt repo, reconfigured with -DCMAKE_PREFIX_PATH=/opt/rocm-7.2.2, rebuilt dflash_server clean. Result: full clean startup, past every point that previously crashed, server reaches listening on http://0.0.0.0:8080, and a real /v1/chat/completions request returns a correct response:

{"choices":[{"finish_reason":"stop","message":{"content":"Hello, how can I help?","role":"assistant"}}],
 "usage":{"completion_tokens":8,"prompt_tokens":12,
   "timings":{"decode_tokens_per_sec":7.2,"prefill_ms":944.8}}}

(Small test prompt, not the full qualification benchmark — this is just confirming genuine end-to-end correctness, not a throughput claim. Haven't yet re-run the full qualification sweep to get a real tok/s comparison against the reported 51.1 tok/s median.)

I don't have a root-cause explanation for why 7.2.2 fixes it (AMD's changelogs between 6.4.1 and 7.2.2 don't call out anything that obviously matches this symptom — closest candidates are a gfx1201 "memory coherency" fix in rocSHMEM 3.2.0, which this build doesn't even use, and a HIP 7.0.2 fix for "a segmentation fault in gpu_metrics ... during CPU synchronization," which is suggestive but not a confirmed match). Given the fix is confirmed regardless of the exact mechanism, might be worth bumping the documented minimum ROCm version for the dual-GPU DS4 path, or at least noting 6.4.1 is known-bad for this configuration.

Happy to share the full build recipe / a Dockerfile if useful for reproducing on your end.

[chrisjohnson]

Filed a follow-up as a separate issue since it's a different bug (in dflash_server's own worker-pool code, not the HIP runtime) found while stress-testing after the fix above: #682 — a heap-use-after-free in Ds4HcMatvecPool (patch included) plus a residual, unresolved rare deadlock in the same code.