EP support for rail isolated topology through Host based RDMA - #558
EP support for rail isolated topology through Host based RDMA#558itej89 wants to merge 134 commits into
Conversation
Add proxy infrastructure for EP on AINIC where GPU-initiated IBGDA WQE posting doesn't work (NIC rejects GPU-written WQEs). The proxy approach: GPU writes command descriptors to a shared ring, CPU thread calls ibv_post_send, GPU polls per-slot status for completion. New files: - proxy_types.hpp: ProxyCmd, ProxyRing shared GPU+CPU types (128-byte aligned) - proxy_device_primitives.hpp: GPU __device__ functions (ProxyPostWrite, ProxyPostWriteInline, ProxyPostAtomic, ProxyQuiet) - proxy_thread.hpp/.cpp: CPU proxy thread (polls ring, ibv_post_send, ibv_poll_cq) - gpu_proxy_rdma_repro.cpp: standalone 2-node reproducer Modified: - internal.hpp: add useProxy + proxyRing to GpuStates - shmem_ibgda_kernels.hpp: add proxy path to ShmemPutMemNbi and ShmemQuietPsd Test results (MI350X + AINIC, loopback): - Level 1 (types): PASS - Level 2 (CPU proxy thread, 1000 ops): PASS, 2.3M ops/s - Level 3 (GPU proxy, 5000 ops with ring wrap): PASS, 107K ops/s, 9.4 us/op Refs: #2 Co-Authored-By: Claude <noreply@anthropic.com>
Wire the proxy infrastructure into the MORI build and runtime: - Add proxy paths to all 6 PSD SHMEM functions in shmem_ibgda_kernels.hpp: ShmemPutMemNbi, ShmemPutSizeImmNbi, ShmemPutMemNbiSignal, ShmemAtomicSizeNonFetch, ShmemAtomicTypeFetch, ShmemQuietPsd - Each function checks globalGpuStates->useProxy and redirects to ProxyPost* primitives, bypassing IBGDA WQE/doorbell/CQ entirely - Allocate ProxyRing in GpuStateInit when MORI_USE_IBGDA_PROXY=1 - Start/shutdown ProxyThread with QP handles from endpoint setup - Add proxy_thread.cpp to application CMakeLists.txt - Forward-declare ProxyThread in internal.hpp (avoids <atomic> in HIP) - Add host-only guard to proxy_thread.hpp Env var: MORI_USE_IBGDA_PROXY=1 enables proxy mode for ionic AINIC. Default (unset or 0) uses native IBGDA — no behavior change. Co-Authored-By: Claude <noreply@anthropic.com>
The ionic IBGDA provider creates ibv_qp via ibv_create_qp_ex but never populated RdmaEndpoint.ibvHandle. The proxy thread needs these handles to call ibv_post_send/ibv_poll_cq. Set endpoint.ibvHandle.qp and .cq from the IonicQpContainer's ibv handles. Also add stderr debug prints to track proxy activation in shmem init. EP test status: proxy activates (32 QPs found, proxy thread started), GPU kernel takes proxy path (no IBGDA WQE posting, no GPU fault), but RDMA fails with remote access error (status=10) — this is a rail topology issue (MR registered on wrong PD/NIC), not a proxy issue. Next: fix rail-isolated topology routing (Task 4). Co-Authored-By: Claude <noreply@anthropic.com>
Standalone test proving any ionic NIC can DMA any GPU's VRAM across XGMI on MI350X. All three scenarios pass: - ionic_0 → GPU 0 (same rail): PASS - ionic_3 → GPU 0 (cross-NIC): PASS - ionic_0 → GPU 5 (reverse cross-NIC): PASS This validates the send-side routing approach: GPU 0 can send to remote GPU 3 by posting through ionic_3, which reads GPU 0's buffer across XGMI and sends to remote ionic_3. Co-Authored-By: Claude <noreply@anthropic.com>
Cherry-pick rail-affinity QP distribution from fix/ep-rail-affinity-ainic: - Add allRdmaDeviceContexts (one per ionic NIC) - Create QP[qp] from allRdmaDeviceContexts[qp % N] - Each QP's GID matches its rail for same-rail routing Add lkey_override to ProxyQpHandle so proxy thread can use per-NIC lkeys when posting through a QP on a different NIC than the buffer's owner. Cross-NIC DMA validated: ionic_3 can DMA GPU 0's VRAM across XGMI (PASS). Still needed: per-NIC MR registration in symmetric memory init. Co-Authored-By: Claude <noreply@anthropic.com>
Add per-NIC MR registration and rkey Allgather for send-side routing: - symmetric_memory.cpp: register heap on all NIC PDs, Allgather rkeys per NIC - init.cpp: wire per-NIC lkeys + rkeys into ProxyQpHandle - proxy_thread: override lkey AND rkey when posting Current status: QP index mapping bug — epIndex from GPU kernel doesn't match the proxy's flat QP array. All posts go to qp_idx=0 instead of the correct per-NIC QP. Transport retry errors from wrong routing. Fix needed: preserve epIndex→QP mapping (include null slots for non-RDMA peers, or add an index translation table). Co-Authored-By: Claude <noreply@anthropic.com>
The GPU kernel passes epIndex (pe * numQpPerPe + qp) as qp_idx to the proxy. The proxy QP array must match this indexing. Changed from a compact push_back array to a full-size indexed array with null entries for non-RDMA peer slots. Current status: SIGSEGV during multi-rail QP creation — ionic's CreateRdmaEndpoint tries to map GPU doorbells via rocm_memory_lock_to_fine_grain when creating QPs on non-affinity NICs. In proxy mode this mapping is unnecessary (CPU calls ibv_post_send, not GPU). Next: skip doorbell mapping in proxy mode, or use plain ibv_create_qp instead of the IBGDA parent domain path. Co-Authored-By: Claude <noreply@anthropic.com>
Status: proxy infrastructure works (Level 1-3 pass), per-NIC MR/rkey exchange works, multi-rail QP creation attempted but SIGSEGV during EP init — likely from IonicQpContainer GPU doorbell mapping when creating QPs on non-affinity NICs. Need to: add debug fprintf traces to find exact crash point, then create plain QPs (no IBGDA parent domain) in proxy mode. Reverted ionic.cpp to clean state to re-apply changes with proper debug tracing rather than guessing. Co-Authored-By: Claude <noreply@anthropic.com>
PYTHONFAULTHANDLER reveals crash at test_dispatch_combine_internode.py:345: self.rng = torch.Generator(device=self.device) The SIGSEGV is in torch.Generator device init AFTER shmem init completes. The GPU's HIP runtime is in a bad state — likely from cross-NIC QP creation touching GPU state on non-affinity devices. In proxy mode, plain QP creation (ibv_create_qp on base PD) avoids GPU VRAM allocation and doorbell mapping. But the IBGDA QPs for the PRIMARY NIC (qpPool path) still go through IonicQpContainer which calls rocm_memory_lock_to_fine_grain and hipGetDevice. The crash varies by process (3, 6, 7) — non-deterministic, suggests a race in GPU device context across the 8 multiprocessing.spawn processes sharing the same GPU set. Need to investigate: does IonicQpContainer corrupt GPU state when called from multiprocessing.spawn child processes? Co-Authored-By: Claude <noreply@anthropic.com>
IonicDeviceContext constructor creates GPU parent domains with custom allocators (hipExtMallocWithFlags) which corrupt GPU state when called from forked multiprocessing.spawn child processes. In proxy mode, skip create_parent_domain entirely — plain QPs don't need GPU VRAM for SQ/CQ. Also skip parent domain for allRdmaDeviceContexts (all 8 NICs). SIGSEGV now occurs at line 343 (shmem_mype()) instead of line 345 (torch.Generator). The crash shifted from GPU device init to MORI shmem global state access. Need to investigate shmem_mype() C extension. Co-Authored-By: Claude <noreply@anthropic.com>
Root cause of SIGSEGV: stale JIT cache (.mori/jit/) compiled with old GpuStates layout (without useProxy/proxyRing fields). hipMemcpy of the larger struct overwrote adjacent GPU globals, corrupting shmem state. Fix: clear JIT cache (rm -rf ~/.mori/jit/) when GpuStates layout changes. The JIT hash should include struct sizes to auto-invalidate, but currently doesn't — needs follow-up. Also: use posix_memalign instead of hipHostMalloc for proxy ring to avoid GPU context corruption in multiprocessing.spawn child processes. hipHostMallocCoherent modifies GPU device state which conflicts with concurrent child processes using the same GPU set. Co-Authored-By: Claude <noreply@anthropic.com>
Re-applied all proxy ionic changes after JIT cache fix: - Skip create_parent_domain in proxy mode - Create plain QPs via ibv_create_qp (no GPU VRAM SQ/CQ) - ConnectEndpoint handles proxyQpPool - Proxy thread starts with per-NIC lkey/rkey overrides Status: no more JIT cache SIGSEGV, but per-NIC ibv_reg_mr in proxy thread start still causes SIGSEGV in some child processes. The ibv_reg_mr on ionic for GPU VRAM from a non-affinity GPU process may corrupt GPU device context. Next: move per-NIC MR registration to symmetric_memory.cpp (which already does ibv_reg_mr safely) instead of doing it in the proxy thread start. Or defer multi-NIC until single-NIC proxy works end-to-end. Co-Authored-By: Claude <noreply@anthropic.com>
…crash - Move per-NIC MR registration from init.cpp to symmetric_memory.cpp where ibv_reg_mr already runs safely - Store perNicLkeys in SymmMemManager, read from init.cpp - Don't start proxy thread when qpCount==0 (no RDMA peers) - Add barrier between per-NIC MR registrations - Add finalize debug traces Confirmed: crash is from proxy thread start (pthread_create + main loop) in multiprocessing.spawn child processes. Without proxy thread, all 8 processes init successfully and reach dispatch (hangs at IBGDA as expected). With proxy thread, random child processes SIGSEGV. Proxy thread issue: ibv_poll_cq in the thread's main loop may access invalid verbs state in forked/spawned child processes. Need to investigate CQ/QP handle validity across process boundaries. Co-Authored-By: Claude <noreply@anthropic.com>
Two fixes that resolve the SIGSEGV in multiprocessing.spawn child processes: 1. Lazy CQ polling: only poll CQ after ops_posted > 0. Previously the thread called ibv_poll_cq immediately on start, before any QP transitions completed across all processes. 2. Use posix_memalign + hipHostRegister instead of hipHostMalloc. hipHostMallocCoherent corrupts GPU device state in child processes. posix_memalign + hipHostRegister(Mapped|Portable) achieves the same GPU accessibility without touching GPU device context at allocation. Result: EP test runs with proxy enabled — no SIGSEGV, all 8 GPU processes init successfully, proxy threads start, GPU kernel writes to proxy ring, CPU thread posts via ibv_post_send, CQE completions received. CQE errors (transport retry) are expected with single NIC due to rail-isolated routing — needs multi-NIC for full EP. Co-Authored-By: Claude <noreply@anthropic.com>
Pensando AINIC has three issues beyond the known IBGDA doorbell rejection: 1. RDMA atomics return CQE success but don't modify remote memory 2. QP connections must use matching NIC on both sides (rail isolation) 3. Traffic class TC must be set on proxy QPs for AINIC routing Fixes applied: - Agreed-rail mapping: both sides use max(myLocalGpu, peerLocalGpu) so QP connections are always on the same NIC index - Proxy QP setup matches non-proxy: active MTU, SL, TC, max_rd_atomic=15, IBV_ACCESS_REMOTE_ATOMIC - Atomic internal buffer (ibuf) allocated and registered per proxy QP for atomic result DMA; atomics use ibuf's own lkey instead of perNic override - IBV_SEND_FENCE on atomics for write-before-atomic ordering - Emulate RDMA atomics via IBV_WR_SEND_WITH_IMM: sender posts inline [dst_addr, add_value] payload; receiver proxy DrainCq detects IBV_WC_RECV, reads payload, does CPU __atomic_fetch_add on GPU memory, re-posts recv WR - Pre-post 128 recv WRs per QP after RTS transition (Pensando rejects post_recv in RESET state) - Recv buffer info plumbed: IonicDeviceContext::proxyRecvInfo map → IBVerbsHandle::recvBuf/recvLkey/recvCount → ProxyQpHandle - NCCL_IB_DISABLE=1 required (NCCL also hits IBGDA doorbell issue) Status: EP dispatch/combine runs end-to-end, 127/128 tokens pass correctness. Token 0 fails due to CPU-GPU atomic coherency race when both intra-node (GPU atomic via P2P) and inter-node (CPU __atomic_fetch_add via proxy) write to the same signal counter. Co-Authored-By: Claude <noreply@anthropic.com>
Two fixes for DispatchInterNodeRecv data corruption on Pensando AINIC: 1. hipDeviceMallocUncached for proxy mode symmetric memory: NIC RDMA writes go directly to GPU VRAM, bypassing GPU L2 cache. Without this, the GPU reads stale L2 cache data while NIC has written fresh data to VRAM, corrupting dispatch metadata (expert indices). 2. PCIe read fence before signal counter atomic: RDMA_WRITE (data) and SEND_WITH_IMM (signal) take different paths to the receiver — data goes NIC→GPU VRAM via DMA, signal goes NIC→CPU recv buffer→proxy thread. A CPU read from GPU VRAM forces PCIe posted write ordering, ensuring the data DMA completes before the signal counter is incremented. Combined with clflush+sfence for CPU→GPU write visibility. Result: Node 1 fully passes (Dispatch + Combine). Node 0 dispatch passes with correct token counts. 126/128 tokens pass data correctness. 2 tokens (token 0 with duplicate PE routing) fail — likely an EP algorithm edge case, not a proxy issue. Co-Authored-By: Claude <noreply@anthropic.com>
128 recv WRs was insufficient for higher-traffic routing patterns, causing combine phase hangs when recv WRs were exhausted before the proxy thread could re-post them. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Proxy pthreads default to GPU 0 — 7/8 threads need hipSetDevice to access the correct GPU's VRAM for __atomic_fetch_add signal counters. Use hipGetDevice to check first, only call hipSetDevice when needed. Co-Authored-By: Claude <noreply@anthropic.com>
hipSetDevice is called once at thread start. The conditional check was unnecessary noise — all proxy pthreads default to GPU 0 and need to be set to their correct GPU. The 2-token data corruption on node 0 is pre-existing — it occurred in early tests before hipSetDevice was introduced. Not caused by the proxy's HIP context setup. Co-Authored-By: Claude <noreply@anthropic.com>
Root cause of 2-token data corruption: SEND_WITH_IMM signal and RDMA_WRITE data travel through different PCIe initiators (CPU vs NIC). PCIe does not order writes from different sources. The GPU sees the CPU-written signal before the NIC's data DMA completes, reads stale staging data, and gets garbage expert indices. Fix: convert PROXY_ATOMIC_FETCH_ADD to IBV_WR_RDMA_WRITE instead of IBV_WR_SEND_WITH_IMM. Both data and signal now go through the NIC, same PCIe initiator. IBV_SEND_FENCE on the same QP guarantees the data write completes before the signal write. No CPU writes to GPU VRAM, no hipSetDevice, no recv WRs needed. This is safe because there is zero contention on signal entries: each receiver GPU's chunkFlag/nodeRecvTokenNum is written by exactly one sender GPU (proxyPe = destNode * gpuPerNode + rank % gpuPerNode, blockFlagCounter is per-GPU with unique flagSlotIds via atomicAdd). Removes: SEND_WITH_IMM recv handling, hipSetDevice, recv WR posting, ProxyRecvInfo, IBVerbsHandle recv fields, proxy uncached heap hack. Net -120 lines. Co-Authored-By: Claude <noreply@anthropic.com>
…_IMM for barriers Two types of atomic operations in EP: 1. Signal writes (chunkFlag, nodeRecvTokenNum) — paired with data, need same NIC→PCIe→VRAM path for ordering. No contention (one sender per entry). New PROXY_SIGNAL_WRITE op → IBV_WR_RDMA_WRITE. 2. Barrier atomics (crossDeviceBarrier) — multiple QPs add to same counter, need real addition. No data ordering requirement. PROXY_ATOMIC_FETCH_ADD → IBV_WR_SEND_WITH_IMM + CPU atomic. GPU kernel uses ProxyPostSignalWrite() for ShmemPutMemNbiSignalThread (data+signal) and ProxyPostAtomicNonFetch() for standalone barrier atomics. Co-Authored-By: Claude <noreply@anthropic.com>
1024 slots was insufficient for 4096+ tokens — each token generates ~2 proxy commands (data write + signal write). With 4096 tokens per GPU, ~8000 commands needed. Ring exhaustion caused the GPU to wait for free slots while the proxy couldn't drain fast enough. 8192 slots handles up to ~4096 tokens per GPU. For bench mode (26K tokens), may need further increase. Co-Authored-By: Claude <noreply@anthropic.com>
65536 slots handles bench mode (~26K tokens, ~50K+ proxy commands). Dispatch passes at all scales. Combine hangs at bench scale due to SEND_WITH_IMM barrier atomics — some CQEs go missing under high load. Test passes: 128, 256, 1024, 4096 tokens (dispatch + combine). Bench (26K tokens): dispatch passes, combine stalls. Co-Authored-By: Claude <noreply@anthropic.com>
IBV_SEND_FENCE on signal writes caused Pensando NIC to stall CQE delivery when many signals queued on the same QP. RC QP guarantees responder-side ordering without FENCE — data RDMA_WRITE completes before signal RDMA_WRITE at the remote GPU. With this fix: - test (128-4096 tokens): Dispatch Pass + Combine Pass, all rounds - bench (26K tokens): Dispatch Pass + Combine Pass, 10 rounds Dispatch: avg 9.24 GB/s RDMA, avg 13.4ms latency Combine: avg 34.56 GB/s RDMA, avg 3.6ms latency Co-Authored-By: Claude <noreply@anthropic.com>
Strip post traces and stall diagnostics for clean bench runs. Co-Authored-By: Claude <noreply@anthropic.com>
Strip all [MoRI-PROXY] debug prints from init, ionic, symmetric_memory. Keep only error prints (CQE error, ibv_post_send failed, RECV CQE error). Co-Authored-By: Claude <noreply@anthropic.com>
The sed-based fprintf removal left orphaned format argument lines that caused compile errors. Clean them up properly. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
RegisterSymmMemObj was doing 8 barriers + 8 Allgathers across all PEs for EVERY shmem_malloc call. With vLLM DeepSeek-V4 (61 MoE layers, ~31 shmem_malloc per layer = ~1900 calls), this meant ~15,000 barriers across 16 PEs, causing MORI SHMEM init to hang for 15+ minutes. Sub-allocations within the StaticHeap share the heap's MR — the per-NIC rkeys from sub-allocation registrations overwrote the heap's keys and were never used by the proxy thread (which reads perNicLkeys from the SymmMemManager, set only by the heap registration). Fix: gate per-NIC MR registration on heap_begin=true. Only the initial heap allocation does the 8-NIC barrier+Allgather dance. Sub-allocations skip it entirely. Co-Authored-By: Claude <noreply@anthropic.com>
- Check wc.status first in DrainCq — on error, opcode may be garbage - Null-check atomic target address from SEND_WITH_IMM payload - Replace fprintf with MORI_LOG_ERROR in proxy_thread.cpp - Replace fprintf with MORI_APP_INFO in init.cpp proxy log line Addresses PR ROCm#558 review comments #1, #3, and #7. Co-Authored-By: Claude <noreply@anthropic.com>
* fix: null-check atomic target address in proxy DrainCq Guard against dereferencing a null pointer from the SEND_WITH_IMM payload. The address comes from the wire — skip the atomic if zero. Addresses PR ROCm#558 review comments #1 and #7 (wc.status check was already implemented in the current code). Co-Authored-By: Claude <noreply@anthropic.com> * fix: check wc.status before opcode, use MORI logger in proxy - Check wc.status first in DrainCq — on error, opcode may be garbage - Null-check atomic target address from SEND_WITH_IMM payload - Replace fprintf with MORI_LOG_ERROR in proxy_thread.cpp - Replace fprintf with MORI_APP_INFO in init.cpp proxy log line Addresses PR ROCm#558 review comments #1, #3, and #7. Co-Authored-By: Claude <noreply@anthropic.com> --------- Co-authored-by: Tej Kiran <kiran.tej@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
When use_native_atomics is true, IBV_WR_ATOMIC_FETCH_AND_ADD DMAs the old value to sge.addr (the ibuf at cmd->src_addr). Read it back into ring->cmds[slot].result before marking PROXY_COMPLETED so the GPU can read the fetch result. No-op for SEND_WITH_IMM path (use_native_atomics=false) — result stays 0 since the old value is on the remote CPU. Co-Authored-By: Claude <noreply@anthropic.com>
There was a problem hiding this comment.
Hi @itej89, Thanks for the refactor!
Could you fix the conflict before merge?
And could you have a test on crusoe cluster again?
And three of my earlier comments look only partly addressed — following up on each rather than re-opening the threads:
1. proxy_thread.cpp:58 — bounds check on the wire pointer
The payload.addr == 0 guard helps, but a null check isn't the property I was asking about. A native IBV_WR_ATOMIC_FETCH_AND_ADD has the responder NIC validate the target address against the MR named by the rkey. A SEND carries no such constraint, so the emulation drops that validation entirely and the CPU does an 8-byte add wherever the payload points. Any wire corruption or a heap-layout mismatch between ranks becomes an arbitrary write into the local process.
GpuStates already tracks heapBaseAddr / heapEndAddr (internal.hpp:128-129); passing those into ProxyThread::Init and rejecting out-of-range payloads would restore the check cheaply. Was there a reason not to?
2. proxy_thread.cpp:140-145 — IBV_SEND_INLINE inline data
This one looks unchanged. PROXY_RDMA_WRITE_INLINE sets IBV_SEND_INLINE but leaves sge.addr as the raw cmd->src_addr assigned at line 122, which is a GPU address — so libibverbs memcpys the payload from a GPU pointer in the proxy thread's host context.
The answer is in the next case down: PROXY_SIGNAL_WRITE (:146-155) correctly redirects sge.addr into the host-side InlineBuf before setting the flag. PROXY_RDMA_WRITE_INLINE needs the same treatment. Its producer is shmem_proxy_kernels.hpp:151.
3. proxy_device_primitives.hpp:158 — fetch atomics returning 0
The fix in DrainCq is gated on qph.use_native_atomics:
if (ring_->cmds[slot].op == PROXY_ATOMIC_FETCH_ADD && qph.use_native_atomics) {
ring_->cmds[slot].result = *reinterpret_cast<volatile uint64_t*>(ring_->cmds[slot].src_addr);
}and use_native_atomics is vendorId == Mellanox (init.cpp:784). So ionic and bnxt — the two NICs this PR is aimed at — still take the SEND_WITH_IMM path and still get 0 back. The emulated path would need the responder to send the pre-add value back, or ProxyPostAtomicFetch should fail loudly rather than return a wrong value silently.
4. New, and I think it should block: proxy_thread.cpp:96 breaks non-x86 builds even with the feature off
asm volatile("clflush (%0)" :: "r"(target) : "memory");
asm volatile("sfence" ::: "memory");There's no #if defined(__x86_64__), and CMakeLists.txt:66 adds proxy_thread.cpp to the source list unconditionally — so an aarch64 build fails whether or not MORI_ENABLE_HOST_PROXY is set. Everything else in this PR is cleanly opt-in (git diff -w shows zero removed lines in init.cpp, and symmetric_memory.cpp's native branch is byte-identical to main), so this is the one place the feature reaches past its own switch.
Guarding the asm is enough; __builtin_ia32_clflush plus std::atomic_thread_fence would also work. Worth noting the clflush is a no-op for its stated purpose anyway if the target is uncached device memory.
clflush/sfence are x86-only and break aarch64 builds. The file is compiled unconditionally. clflush is also a no-op on uncached device memory (hipDeviceMallocUncached). Replace with portable std::atomic_thread_fence(std::memory_order_seq_cst). Addresses PR ROCm#558 round-2 review comment R2-4. Co-authored-by: Tej Kiran <kiran.tej@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
PROXY_RDMA_WRITE_INLINE set IBV_SEND_INLINE but left sge.addr as the GPU VRAM address from cmd->src_addr. libibverbs memcpys from sge.addr at post time — a GPU pointer the CPU cannot access. Fix: add inline_data[48], inline_tag, and inline_len fields to ProxyCmd (uses existing padding, struct stays 128 bytes). GPU copies data into cmd->inline_data via 8-byte stores on host-pinned ring, sets tag to PROXY_INLINE_IMM_WRITE and len. CPU validates tag+len before reading. - proxy_types.hpp: add ProxyInlineTag enum, inline_data[48], inline_tag, inline_len fields, PROXY_MAX_INLINE_DATA constant, static_assert - proxy_device_primitives.hpp: ProxyPostWriteInline copies src into cmd->inline_data, sets tag and len - proxy_thread.cpp: BuildWr validates tag+len, reads from cmd->inline_data - shmem_proxy_kernels.hpp: pass val pointer directly (not cast to u64) Tagged payload design: inline_tag tells CPU what inline_data contains. Future callers add new tags for new payload types without protocol changes. Addresses PR ROCm#558 round-2 review comment R2-2. Co-Authored-By: Claude <noreply@anthropic.com>
PROXY_RDMA_WRITE_INLINE set IBV_SEND_INLINE but left sge.addr as the GPU VRAM address from cmd->src_addr. libibverbs memcpys from sge.addr at post time — a GPU pointer the CPU cannot access. Fix: add inline_data[48], inline_tag, and inline_len fields to ProxyCmd (uses existing padding, struct stays 128 bytes). GPU copies data into cmd->inline_data via 8-byte stores on host-pinned ring, sets tag to PROXY_INLINE_IMM_WRITE and len. CPU validates tag+len before reading. - proxy_types.hpp: add ProxyInlineTag enum, inline_data[48], inline_tag, inline_len fields, PROXY_MAX_INLINE_DATA constant, static_assert - proxy_device_primitives.hpp: ProxyPostWriteInline copies src into cmd->inline_data, sets tag and len - proxy_thread.cpp: BuildWr validates tag+len, reads from cmd->inline_data - shmem_proxy_kernels.hpp: pass val pointer directly (not cast to u64) Tagged payload design: inline_tag tells CPU what inline_data contains. Future callers add new tags for new payload types without protocol changes. Addresses PR ROCm#558 round-2 review comment R2-2. Co-Authored-By: Claude <noreply@anthropic.com>
PROXY_RDMA_WRITE_INLINE set IBV_SEND_INLINE but left sge.addr as the GPU VRAM address from cmd->src_addr. libibverbs memcpys from sge.addr at post time — a GPU pointer the CPU cannot access. Fix: add inline_data[48], inline_tag, and inline_len fields to ProxyCmd (uses existing padding, struct stays 128 bytes). GPU copies data into cmd->inline_data via 8-byte stores on host-pinned ring, sets tag to PROXY_INLINE_IMM_WRITE and len. CPU validates tag+len before reading. - proxy_types.hpp: add ProxyInlineTag enum, inline_data[48], inline_tag, inline_len fields, PROXY_MAX_INLINE_DATA constant, static_assert - proxy_device_primitives.hpp: ProxyPostWriteInline copies src into cmd->inline_data, sets tag and len - proxy_thread.cpp: BuildWr validates tag+len, reads from cmd->inline_data - shmem_proxy_kernels.hpp: pass val pointer directly (not cast to u64) Tagged payload design: inline_tag tells CPU what inline_data contains. Future callers add new tags for new payload types without protocol changes. Addresses PR ROCm#558 round-2 review comment R2-2. Co-Authored-By: Claude <noreply@anthropic.com>
PROXY_RDMA_WRITE_INLINE set IBV_SEND_INLINE but left sge.addr as the GPU VRAM address from cmd->src_addr. libibverbs memcpys from sge.addr at post time — a GPU pointer the CPU cannot access. Fix: add inline_data[48], inline_tag, and inline_len fields to ProxyCmd (uses existing padding, struct stays 128 bytes). GPU copies data into cmd->inline_data via 8-byte stores on host-pinned ring, sets tag to PROXY_INLINE_IMM_WRITE and len. CPU validates tag+len before reading. - proxy_types.hpp: add ProxyInlineTag enum, inline_data[48], inline_tag, inline_len fields, PROXY_MAX_INLINE_DATA constant, static_assert - proxy_device_primitives.hpp: ProxyPostWriteInline copies src into cmd->inline_data, sets tag and len - proxy_thread.cpp: BuildWr validates tag+len, reads from cmd->inline_data - shmem_proxy_kernels.hpp: pass val pointer directly (not cast to u64) Tagged payload design: inline_tag tells CPU what inline_data contains. Future callers add new tags for new payload types without protocol changes. Addresses PR ROCm#558 round-2 review comment R2-2. Co-Authored-By: Claude <noreply@anthropic.com>
PROXY_RDMA_WRITE_INLINE set IBV_SEND_INLINE but left sge.addr as the GPU VRAM address from cmd->src_addr. libibverbs memcpys from sge.addr at post time — a GPU pointer the CPU cannot access. Fix: add inline_data[48], inline_tag, and inline_len fields to ProxyCmd (uses existing padding, struct stays 128 bytes). GPU copies data into cmd->inline_data via 8-byte stores on host-pinned ring, sets tag to PROXY_INLINE_IMM_WRITE and len. CPU validates tag+len before reading. - proxy_types.hpp: add ProxyInlineTag enum, inline_data[48], inline_tag, inline_len fields, PROXY_MAX_INLINE_DATA constant, static_assert - proxy_device_primitives.hpp: ProxyPostWriteInline copies src into cmd->inline_data, sets tag and len - proxy_thread.cpp: BuildWr validates tag+len, reads from cmd->inline_data - shmem_proxy_kernels.hpp: pass val pointer directly (not cast to u64) Tagged payload design: inline_tag tells CPU what inline_data contains. Future callers add new tags for new payload types without protocol changes. Addresses PR ROCm#558 round-2 review comment R2-2. Co-authored-by: Tej Kiran <kiran.tej@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
…23) Native IBV_WR_ATOMIC_FETCH_AND_ADD validates the target address via the MR rkey. The SEND_WITH_IMM emulation drops that — the CPU does an atomic wherever the payload points, so wire corruption becomes an arbitrary write. Pass heapBaseAddr/heapEndAddr from GpuStates into ProxyThread::Init and reject payload.addr outside [heapBase, heapEnd) in DrainCq. Addresses PR ROCm#558 round-2 review comment R2-1. Co-authored-by: Tej Kiran <kiran.tej@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
ProxyPostAtomicFetch sets flags=2 ("fetch result required"). CPU
DrainCq checks this flag: if use_native_atomics, writes result from
ibuf; if emulated (SEND_WITH_IMM) and fetch required, logs error
and returns PROXY_ERROR instead of silent wrong value.
Native atomics (IBV_WR_ATOMIC_FETCH_AND_ADD) remain Mellanox-only.
Tested on DO AINIC with ionic — CQE errors confirm ionic hardware
does not execute RDMA atomics correctly (transport retry exceeded).
Addresses PR ROCm#558 round-2 review comment R2-3.
Co-Authored-By: Claude <noreply@anthropic.com>
Add ProxyCmdFlags enum (PROXY_FLAGS_DEFAULT, PROXY_FLAGS_FETCH_REQUIRED). ProxyPostAtomicFetch sets PROXY_FLAGS_FETCH_REQUIRED. CPU DrainCq: - Any non-default flag on emulated path → MORI_APP_ERROR + PROXY_ERROR - PROXY_FLAGS_FETCH_REQUIRED on native path → writes ibuf result Native atomics remain Mellanox-only. Confirmed ionic native atomics fail (CQE transport retry exceeded on DO AINIC). Addresses PR ROCm#558 round-2 review comment R2-3. Co-Authored-By: Claude <noreply@anthropic.com>
Add ProxyCmdFlags enum (PROXY_FLAGS_DEFAULT, PROXY_FLAGS_FETCH_REQUIRED). ProxyPostAtomicFetch sets PROXY_FLAGS_FETCH_REQUIRED. CPU DrainCq: - Any non-default flag on emulated path → MORI_APP_ERROR + PROXY_ERROR - PROXY_FLAGS_FETCH_REQUIRED on native path → writes ibuf result Native atomics remain Mellanox-only. Confirmed ionic native atomics fail (CQE transport retry exceeded on DO AINIC). Addresses PR ROCm#558 round-2 review comment R2-3. Co-Authored-By: Claude <noreply@anthropic.com>
Implement full fetch-and-add return value on emulated path: - ProxyPostAtomicFetch sets PROXY_FLAGS_FETCH_REQUIRED + slot in inline_data - Sender SEND includes [dst_addr, val, qp_idx, slot] (32 bytes, imm=0xA70F) - Receiver does __atomic_fetch_add, sends reply [old_val, slot] (imm=0xA71C) - Sender recv CQE matches reply, writes result to ring, marks COMPLETED - Sender send CQE skips COMPLETED for fetch-required (reply does it) - GPU asserts on PROXY_ERROR if emulated path can't fulfill fetch ProxyPostAtomicNonFetch unchanged (PROXY_FLAGS_DEFAULT, fire-and-forget). Native atomics (CX7) unchanged. Verified: forced all atomics through round-trip on DO AINIC — Dispatch Pass, Combine Pass, ~18% combine overhead vs fire-and-forget. Normal path (NonFetch) has zero overhead. EXPERIMENTAL — not for merge. Co-authored-by: Tej Kiran <kiran.tej@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
Hi @jhchouuu, thank you for the feedback. I've made the requested changes as follows:
|
Integrate proxy RDMA path with upstream's rail-only refactor: - BuildInitialEndpoints replaces BuildAndConnectInitialEndpoints - Proxy QP creation uses GetRailContext within shouldConnect() gating - Both proxyEnabled and railOnly members coexist in Context Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
|
cc: @jhchouuu EP-Over-HOST-PROXY Verification Report - PR(1781fa1) vs mainBranch: ResultsDispatch Performance (bfloat16)
Combine Performance (bfloat16)
Analysis
Cluster Details
|
jhchouuu
left a comment
There was a problem hiding this comment.
LGTM overall — thanks for working through all of these. Just the one last item above to sort out.
One other thing: async_ll still doesn't seem to work under the proxy — it faults reproducibly right after the first dispatch. v1 and v1_ll are both fine. Is that expected / out of scope for this PR, or worth a look before merging?
| ibv_send_wr rwr{}, *rbad = nullptr; | ||
| rwr.opcode = IBV_WR_SEND_WITH_IMM; | ||
| rwr.imm_data = htonl(PROXY_IMM_ATOMIC_REPLY); | ||
| rwr.send_flags = IBV_SEND_SIGNALED | IBV_SEND_INLINE; | ||
| rwr.sg_list = &rsge; |
There was a problem hiding this comment.
@itej89 Would you please take a look at this items?
Blocking: the atomic-fetch reply WR has no wr_id, so its completion corrupts ring slot 0
160: ibv_send_wr rwr{}, *rbad = nullptr; // wr_id zero-initialised
163: rwr.send_flags = IBV_SEND_SIGNALED | IBV_SEND_INLINE; // always produces a send CQE
166: ibv_post_send(qph.qp, &rwr, &rbad); // wr_id never assignedsend_cq == recv_cq (ibverbs.cpp:213-214), so this completion lands in the CQ DrainCq polls. Opcode is IBV_WC_SEND, so it skips the RECV branch and hits the tail:
185: uint32_t slot = static_cast<uint32_t>(wc[i].wr_id) & PROXY_RING_MASK; // = 0
195: ring_->cmds[slot].status = PROXY_COMPLETED;So every atomic-fetch reply marks slot 0 of the responder's own ring as completed. Worst case isn't a lost completion — if slot 0 is PENDING but not yet consumed, MainLoop's if (cmd->status != PROXY_PENDING) break; fires without advancing next_slot_, so that ring stops posting permanently while the GPU's ProxyQuiet sees COMPLETED and moves on. Same route at :106 would mark slot 0 PROXY_ERROR.
Fix — high-bit sentinel, tested before the mask (~0ull & PROXY_RING_MASK is 65535, a valid slot):
static constexpr uint64_t PROXY_WRID_INTERNAL = 1ull << 63; // proxy_types.hpp
rwr.wr_id = PROXY_WRID_INTERNAL; // at :160
if (wc[i].wr_id & PROXY_WRID_INTERNAL) continue; // before :106 and :185Testing: no EP operator uses fetch atomics (dispatch/combine use ShmemAtomicTypeNonFetchThread / ShmemPutMemNbiSignalThread), so the benchmarks can't have covered this path.
…corruption (#26) * fix: tag atomic-fetch reply WR with sentinel wr_id to prevent slot 0 corruption The reply SEND_WITH_IMM in the atomic-fetch round-trip path had wr_id=0 (zero-initialized). Since send_cq == recv_cq and the WR is SIGNALED, its completion landed in DrainCq as IBV_WC_SEND with wr_id=0, which unconditionally marked ring slot 0 as COMPLETED — corrupting whatever command the GPU had pending in that slot. Fix: set wr_id = PROXY_WRID_INTERNAL (high-bit sentinel) on the reply WR and skip any CQE with that sentinel at the top of DrainCq. Co-Authored-By: Claude <noreply@anthropic.com> * fix: log error when internal reply SEND fails Don't silently drop failed reply completions — log the status so operators can diagnose fetch-atomic hangs. Co-Authored-By: Claude <noreply@anthropic.com> * docs: clarify why bit 63 is collision-free for PROXY_WRID_INTERNAL Co-Authored-By: Claude <noreply@anthropic.com> --------- Co-authored-by: Tej Kiran <kiran.tej@amd.com> Co-authored-by: Claude <noreply@anthropic.com>
Motivation
Enable Expert Parallelism (EP) on clusters where GPU-initiated IBGDA doorbells don't work or the vendor DV library is unavailable.
Technical Details
Adds TransportType::PROXY = 3 to the existing dispatch framework. EP kernels route through DISPATCH_TRANSPORT_TYPE → template specializations in shmem_proxy_kernels.hpp, leaving shmem_ibgda_kernels.hpp untouched. Proxy code compiles only when -DMORI_PROXY_ENABLED is set through JIT pipeline.
Per-NIC CPU threads poll host-pinned ring buffers (64K slots × 128B, hipHostRegister mapped). GPU warps write ProxyCmd descriptors; CPU threads convert to ibv_post_send with per-NIC lkey/rkey override. Signal delivery uses RDMA_WRITE for data signals and SEND_WITH_IMM + CPU-side __atomic_fetch_add for barrier atomics (ionic/bnxt). CX7 uses native IBV_WR_ATOMIC_FETCH_AND_ADD for barrier atomics.
Supported NIC backends: Pensando AINIC (ionic), Mellanox CX7 (mlx5), Broadcom Thor2 (bnxt).
When
MORI_EP_OVER_RDMA=1, the DV API availability check is skipped for all providers — proxy uses plain ibverbs and does not require vendor DV libraries.How to Run
Test Results
Commit:
9ae6ebdf— 2-node EP16 benchmarkProxy mode (MORI_EP_OVER_RDMA=1)
Native IBGDA (no proxy, CX7 only)
AINIC native: IBGDA hangs (GPU-written WQEs not visible to NIC on VFIO platform).
Broadcom native: DV library not available on test cluster (RDMA uABI 1).
Clusters tested
Submission Checklist