[Draft] EP support for rail isolated topology through Host based RDMA - #558
Draft
itej89 wants to merge 110 commits into
Draft
[Draft] EP support for rail isolated topology through Host based RDMA#558itej89 wants to merge 110 commits into
itej89 wants to merge 110 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>
Co-Authored-By: Claude <noreply@anthropic.com>
Warp-scope and Block-scope proxy kernel specializations were calling the Thread-scope implementation from ALL lanes/threads, causing 64x duplicate proxy ring posts per warp-scope call. Match v3's behavior: WarpKernel gates on laneId==0, BlockKernel gates on FlatBlockThreadId()==0. Co-Authored-By: Claude <noreply@anthropic.com>
This reverts commit 3171441.
Match v3's WarpKernelImpl/BlockKernelImpl pattern: only lane 0 (Warp) or thread 0 (Block) enters the Thread-scope proxy post. Without this, all 64 lanes post duplicate commands — 64x more SEND_WITH_IMM atomics flooding remote proxy threads with CQE draining, causing 20-40x slowdown. Co-Authored-By: Claude <noreply@anthropic.com>
Replace max(myGpu, peerGpu) % N with (myGpu + peerGpu) % N. max() is asymmetric: rank 7 funnels all traffic through NIC 7 (1 thread), while rank 0 uses all 8 NICs (7 threads). Addition is commutative (both sides agree) and distributes peers evenly — every rank uses 7 of 8 NICs. Co-Authored-By: Claude <noreply@anthropic.com>
This reverts commit 7e010ad.
Remove [MoRI-DBG] fprintf traces from init.cpp, runtime.cpp, proxy_thread.cpp, symmetric_memory.cpp. Keep CQE error prints (useful for production debugging). Restore test script to original. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Add useProxy, proxyRings[], proxyQuietHead[], numProxyRings, numNics, localGpuIdx back into GpuStates alongside ProxyGpuStates. First step to match v3 behavior for contention investigation. Co-Authored-By: Claude <noreply@anthropic.com>
Proxy kernels now read rings/quietHead/numNics/localGpuIdx from GpuStates directly instead of ProxyGpuStates. init.cpp populates both structs. This makes proxy fields travel via the same hipMemcpy as GpuStates (no separate symbol lookup). Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Remove ProxyGpuStates struct, globalProxyState device symbol, separate hipMemcpy path, and all references. Proxy fields (rings, quietHead, numProxyRings, numNics, localGpuIdx) are in GpuStates and travel via the single existing GpuStates hipMemcpy. This eliminates the two-symbol contention issue that caused uneven per-rank bandwidth under NIC load. Also simplify ShmemQuietThread() no-arg to use numProxyRings > 0 check instead of PE loop (single device memory read vs 16). Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Reference old GpuStates fields that no longer exist. Co-Authored-By: Claude <noreply@anthropic.com>
Add proxy QP creation and connection to Mlx5DeviceContext, mirroring the existing IonicDeviceContext proxy path. When MORI_EP_OVER_RDMA=1: - CreateRdmaEndpoint uses ibv_create_cq + ibv_create_qp (plain verbs, not DevX) and stores in proxyQpPool - ConnectEndpoint uses ibv_modify_qp RST→INIT→RTR→RTS and pre-posts 512 recv WRs for SEND_WITH_IMM atomic emulation - context.cpp dynamic_cast handles both IonicDeviceContext and Mlx5DeviceContext for GetProxyRecvInfo CX7 on MI300x has 8 mlx5 NICs — same topology as AINIC (8 proxy threads, 8 rings, agreedRail spreads QPs across all NICs). Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Traces in proxy_thread.cpp (commands posted, errors, null QP) and init.cpp (per-NIC QP count, total threads). Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
…during hang 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>
CX7 supports IBV_WR_ATOMIC_FETCH_AND_ADD — NIC hardware does the atomic directly to GPU VRAM. Ionic doesn't, so it keeps SEND_WITH_IMM emulation where CPU does __atomic_fetch_add. Add use_native_atomics flag to ProxyQpHandle, set true for Mellanox vendor, false for Pensando. BuildWr checks the flag per QP. 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>
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.
Motivation
Enable Expert Parallelism (EP) on rail-isolated clusters where Cross NIC communication isn't possible and GPU-initiated IBGDA doorbells won't work.
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 for rail-isolated MR registration. Signal delivery uses RDMA_WRITE for data signals and SEND_WITH_IMM + CPU-side __atomic_fetch_add for barrier atomics (ionic). CX7 uses native IBV_WR_ATOMIC_FETCH_AND_ADD for barrier atomics.
Test Plan
MORI EP benchmark tests
Test Result
TBD
Submission Checklist