EP over RDMA v6: TransportType::PROXY — native dispatch framework - #5
Open
itej89 wants to merge 88 commits into
Open
EP over RDMA v6: TransportType::PROXY — native dispatch framework#5itej89 wants to merge 88 commits into
itej89 wants to merge 88 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>
Consistent with the type name ProxyGpuStates. init.cpp references need updating (cpp pass). Co-Authored-By: Claude <noreply@anthropic.com>
- init.cpp: rename proxyGpuState → proxyGpuStates (16 places) - runtime.cpp: restore two-symbol copy — GpuStates always copied via globalGpuStates, ProxyGpuStates copied via globalProxyState when proxy enabled (using commContext->IsProxyEnabled()) Co-Authored-By: Claude <noreply@anthropic.com>
Host uses commContext->IsProxyEnabled(), device uses transportTypes[pe]. No need for a redundant active flag. Co-Authored-By: Claude <noreply@anthropic.com>
Remove separate useProxy bool. Use proxyEnabled member directly. Proxy path is a clean if block — no mixed ternary/conditional with native path. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Keep main's original code in else path untouched. Proxy path in if(proxyEnabled) block — one block per operation, no mixed ctx variable. Co-Authored-By: Claude <noreply@anthropic.com>
Add Context::GetRailContext(peerRank) that encapsulates the max(localGpu, peerLocalGpu) % numNics agreed-rail formula. Replaces 4 inline copies in context.cpp. Co-Authored-By: Claude <noreply@anthropic.com>
…led() Remove regCount debug logging. Use context.IsProxyEnabled() directly. Original MR registration path untouched in else block. Co-Authored-By: Claude <noreply@anthropic.com>
Main's original code (ibv_reg_mr + Allgather) left as-is. All proxy additions (heap key caching + per-NIC MR) under one if(IsProxyEnabled()) block after. Co-Authored-By: Claude <noreply@anthropic.com>
Proxy heap: register + allgather + cache + per-NIC MR. Proxy sub-alloc: reuse cached keys, skip allgather (matches v3). Native: main's original code in else block, untouched. Co-Authored-By: Claude <noreply@anthropic.com>
Read env var once, use cached member for all 3 checks. Co-Authored-By: Claude <noreply@anthropic.com>
First proxyEnabled block returns — second was unreachable. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Keep commented-out cqPool.insert + qpPool.insert in their original position from main. Co-Authored-By: Claude <noreply@anthropic.com>
Only proxy QPs need ibvHandle.qp/cq (already set in proxy block). Native IBGDA path doesn't use them — matches main. Co-Authored-By: Claude <noreply@anthropic.com>
GpuStates always defined. Proxy state conditionally appended via _MORI_DEFINE_PROXY_STATE. No duplication. Co-Authored-By: Claude <noreply@anthropic.com>
PutMemNbi: add chunking loop with VmmQueryLocalKey/VmmQueryRemoteAddr for VMM heap, matching v3's behavior. PutSizeImmNbi: add VmmLookupRemote for VMM heap addressing. Static heap path unchanged. 1:1 match with v3 for both heap modes. Co-Authored-By: Claude <noreply@anthropic.com>
Matches the constant name in internal.hpp after shmem_proxy_state.hpp was deleted. 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>
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.
2 tasks
itej89
added a commit
that referenced
this pull request
Aug 27, 2026
…path (#18) 1. Restore rdmaRegister check in proxy path (line 198) — without it, SDMA-only buffers that shouldn't be registered hit ibv_reg_mr, which can fail on ionic's single-MR limit. 2. Guard heapRkeys_ memcpy against empty vector — in ShmemMode::Isolation the heap is never registered so heapRkeys_ is empty, causing memcpy from nullptr. Addresses PR ROCm#558 review comments #5 and #6. Co-authored-by: Tej Kiran <kiran.tej@amd.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.
Summary
TransportType::PROXY = 3to the existing transport dispatch frameworkShmemPutMemNbiThreadKernel<PROXY>) — same pattern as RDMA/P2P/SDMADISPATCH_TRANSPORT_TYPEmacro — no custom guard macros neededMORI_EP_OVER_RDMA=1, init setstransportTypes[pe] = PROXYfor inter-node RDMA peersChanges to shmem_device_api.hpp vs main
shmem_proxy_kernels.hpp)else if (PROXY)branches in dispatch macrosShmemQuietThread()no-arg: PROXY check before RDMA fallback#ifdefguardsArchitecture (vs v5)
v5 used
PROXY_DISPATCH_GUARDmacros outside the framework. v6 uses the framework's ownDISPATCH_TRANSPORT_TYPE— exactly how SDMA and P2P are dispatched.Test plan
Generated with Claude Code