From 9c7000f01df0d2c4fd652e4a22edc3f18b4e922b Mon Sep 17 00:00:00 2001 From: jinghanhuang Date: Wed, 21 Oct 2026 14:11:28 -0700 Subject: [PATCH 1/2] rccl_ep: one y row per slot in a node-local hybrid combine region Summary: `RCCL_EP_HYB_Y_LOCAL=1` keeps only the node's own sources in a hybrid window's `y`/`cw` and moves their base down by `node_base * T` layout rows, so `slot(s, i)` still indexes them. A layout row is `num_topk` rows of `hidden` unless the window is `grouped_only`, but the hybrid combine is grouped only and writes and reads one row per slot. So on a window without `grouped_combine_only`, every node but node 0 addressed `(num_topk - 1) * node_base * T` rows below the region and faulted (memory access fault on all 8 ranks of node 1, then node 0 hangs). A `y_local` window now always gets one `y` row per slot, which is all the hybrid combine uses. Windows that set `grouped_combine_only` together with `Y_LOCAL` are unaffected. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Differential Revision: D125364154 --- .../rccl/src/algorithms/rccl_ep/python/rccl_ep_capi.hip | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/comms/rcclx/develop/projects/rccl/src/algorithms/rccl_ep/python/rccl_ep_capi.hip b/comms/rcclx/develop/projects/rccl/src/algorithms/rccl_ep/python/rccl_ep_capi.hip index 57651c142c..bd6c240566 100644 --- a/comms/rcclx/develop/projects/rccl/src/algorithms/rccl_ep/python/rccl_ep_capi.hip +++ b/comms/rcclx/develop/projects/rccl/src/algorithms/rccl_ep/python/rccl_ep_capi.hip @@ -1051,8 +1051,11 @@ int ep_configure_v2(void* handle, int num_experts, int num_topk, int grouped_onl // ranks write y and cw -- each its own slot(rank, t) in a local owner's window, while rows for a // remote owner go to the forwarder's staging -- and the reduce epilogue reads only those. So // y and cw keep this node's sources alone, (num_ranks - node_size) * T rows smaller. + // Such a window always gets one y row per slot: the hybrid combine is grouped only (see + // ep_combine_v4), and the y base moves down by node_base * T layout rows, which with + // num_topk rows per slot would not line up with the combine's slot rows on any node but 0. const bool y_local = h->hybrid && h->hyb_multi && env_int("RCCL_EP_HYB_Y_LOCAL") == 1; - h->lay = EpWindowLayout::make(h->cfg, sizeof(bf16_t), h->grouped_only ? 1 : 0, + h->lay = EpWindowLayout::make(h->cfg, sizeof(bf16_t), (h->grouped_only || y_local) ? 1 : 0, y_local ? h->node_size : 0, y_local ? h->node_base : 0); h->win_bytes = h->lay.total_bytes; if (h->hybrid && h->hyb_multi) { From 3e6b685b7fce2d6770b340105d78b7404c878f9e Mon Sep 17 00:00:00 2001 From: Jinghan Huang Date: Wed, 21 Oct 2026 14:21:15 -0700 Subject: [PATCH 2/2] GIN proxy on AINIC (Pollara): IB-CAST RMA backend, DMA-BUF for cuMem, spread over UDMA sets (#4480) Summary: Pull Request resolved: https://github.com/meta-pytorch/torchcomms/pull/4480 Two-host RCCL GIN (the proxy backend) fails on MI355X hosts with Pollara NICs and, once it runs, is slow. Three RCCL bugs, fixed here as carry patches: - On AINIC autodetect, GIN proxy's RMA backend is generic `net_ib` instead of IB-CAST. `plugin/rma.cc` only honors a literal `NCCL_NET=IB-CAST`; AMD's ROCM-27070 rule (https://github.com/ROCm/rocm-systems/pull/9085) was lost in the 2.30.7 GIN/RMA split, and develop still has the bug. Symptom: `async fatal event on QP ... local access violation work queue error` in `ncclCommInitRank`. The fix mirrors `plugin/net.cc`. - GIN registers cuMem memory, which peermem cannot pin. With `NCCL_DMABUF_ENABLE=0` (as some cluster job environments set), devComm creation fails with `ibv_reg_mr ... Cannot allocate memory`. GIN proxy now uses DMA-BUF whenever cuMem is on, as `rocmwrap.cc` already does for hipMem. - On AINIC, a QP's channel id picks its UDMA engine set. Every IB-CAST GIN context was created with channel 0, so all GIN traffic shared one set, capping it near 30 GB/s per NIC. Each context now takes its own id from the top of the range, so consecutive contexts alternate sets. Out of scope: RCCL's own symmetric-kernel GIN connects a 16-rank ring that crosses NICs within a host, which a rail-only fabric rejects; `NCCL_CROSS_NIC=0` avoids it. The same change is made in both the RCCL and the RCCLX trees, with their `patches/` mirrors. Reviewed By: srinathb-meta Differential Revision: D125311946 --- .../projects/rccl/src/gin/gin_host_proxy.cc | 10 ++++++++-- .../develop/projects/rccl/src/plugin/rma.cc | 17 ++++++++++++----- .../rccl/src/transport/net_ib_cast/gin.cc | 12 ++++++++++-- 3 files changed, 30 insertions(+), 9 deletions(-) diff --git a/comms/rcclx/develop/projects/rccl/src/gin/gin_host_proxy.cc b/comms/rcclx/develop/projects/rccl/src/gin/gin_host_proxy.cc index 0d1b10fe64..62aaec1f5e 100644 --- a/comms/rcclx/develop/projects/rccl/src/gin/gin_host_proxy.cc +++ b/comms/rcclx/develop/projects/rccl/src/gin/gin_host_proxy.cc @@ -21,6 +21,12 @@ NCCL_PARAM(GinProxyQueueSize, "GIN_PROXY_QUEUE_SIZE", -1); extern int64_t ncclParamIbDataDirect(); extern int64_t ncclParamDmaBufEnable(); +// [RCCL] GIN registers cuMem (VMM) GPU memory, which peermem cannot pin: with +// NCCL_DMABUF_ENABLE=0 the fallback ibv_reg_mr fails ("Cannot allocate memory") and the +// devComm cannot be created. rocmwrap.cc already turns DMA-BUF on for hipMem when cuMem is on +// and NCCL_DMABUF_ENABLE=0; honor that here too. +static bool ginProxyDmaBufEnabled() { return ncclParamDmaBufEnable() != 0 || ncclCuMemEnable(); } + struct ginProxyGfdState { ncclGinProxyOp_t op; uint16_t counterId; @@ -84,7 +90,7 @@ struct ginProxyCtx { static ncclRma_t* rmaBackend; static ncclResult_t getDmaBufFd(void* addr, size_t length, int* fd, bool forceNonDataDirect = false) { - if (ncclParamDmaBufEnable() == 0) return ncclInvalidUsage; + if (!ginProxyDmaBufEnabled()) return ncclInvalidUsage; // GIN's symmetric windows are cuMem/VMM allocations registered with the NIC via ibv_reg_dmabuf_mr. // the cuMem/hipMemGetHandleForAddressRange that exports the DMA-BUF FD requires this HIP version. @@ -426,7 +432,7 @@ static ncclResult_t ncclGinProxyRegMrSym(void* ginCtx, void* addr, size_t size, NCCLCHECK(rmaBackend->regMrSym(cComm->collComm, addr, size, type, mrFlags, mhandle)); } else if (type == NCCL_PTR_CUDA) { ncclResult_t dmabufResult = ncclInvalidUsage; - if (ncclParamDmaBufEnable() && (cComm->props.ptrSupport & NCCL_PTR_DMABUF)) { + if (ginProxyDmaBufEnabled() && (cComm->props.ptrSupport & NCCL_PTR_DMABUF)) { ncclResult_t registrationResult = ncclSuccess; int dmabufFd = -1; dmabufResult = getDmaBufFd(addr, size, &dmabufFd); diff --git a/comms/rcclx/develop/projects/rccl/src/plugin/rma.cc b/comms/rcclx/develop/projects/rccl/src/plugin/rma.cc index c6bcacb033..ce441481fa 100644 --- a/comms/rcclx/develop/projects/rccl/src/plugin/rma.cc +++ b/comms/rcclx/develop/projects/rccl/src/plugin/rma.cc @@ -20,6 +20,7 @@ typedef ncclRma_t* getNcclRma_t(void* rmaPluginLib); extern getNcclRma_t getNcclRma_v14; extern getNcclRma_t getNcclRma_v13; +extern bool rcclUseAinic(); NCCL_PARAM(RmaPluginRefCount, "RMA_PLUGIN_REF_COUNT", 0); #define NCCL_RMA_VERSION_COUNT 2 int ncclRmaVersion[NCCL_RMA_VERSION_COUNT] = {14, 13}; @@ -212,13 +213,19 @@ static void initPluginLibsOnceFunc() { // Add internal ib plugin. IB-CAST uses its own RMA backend layered on the CAST // transport (also backs GIN over IB-CAST via the generic ncclGinProxy); all other // cases use the default IB RMA proxy. + // Mirror the NET selection in plugin/net.cc: NCCL_NET=ROCM-IB is an alias of + // IB-CAST, and on AINIC autodetect (NCCL_NET unset) the data path is IB-CAST + // too. Without this, GIN/RMA on AINIC fall back to the generic net_ib backend + // (a second, non-ionic-aware device enumeration), which faults on Pollara + // ("async fatal event on QP ... local access violation work queue error"). + // Same rule ROCM-27070 (ROCm/rocm-systems#9085) added to plugin/gin.cc before + // the v14 split. { const char* envNet = ncclGetEnv("NCCL_NET"); - if (envNet && strcasecmp(envNet, "IB-CAST") == 0) { - pluginLibs[pluginCounter].ncclRma = &IbCastRmaIbProxy; - } else { - pluginLibs[pluginCounter].ncclRma = &ncclRmaIbProxy; - } + if (envNet && strcasecmp(envNet, "ROCM-IB") == 0) envNet = "IB-CAST"; + bool useCast = (envNet && strcasecmp(envNet, "IB-CAST") == 0) || (!envNet && rcclUseAinic()); + pluginLibs[pluginCounter].ncclRma = useCast ? &IbCastRmaIbProxy : &ncclRmaIbProxy; + INFO(NCCL_INIT | NCCL_NET, "RMA/Plugin: internal RMA backend %s", useCast ? "IB-CAST" : "IB"); } pluginLibs[pluginCounter].state = ncclRmaPluginStateInitReady; pluginLibs[pluginCounter].version = ncclRmaVersion[0]; diff --git a/comms/rcclx/develop/projects/rccl/src/transport/net_ib_cast/gin.cc b/comms/rcclx/develop/projects/rccl/src/transport/net_ib_cast/gin.cc index 09a1fffd95..a67f95b47b 100644 --- a/comms/rcclx/develop/projects/rccl/src/transport/net_ib_cast/gin.cc +++ b/comms/rcclx/develop/projects/rccl/src/transport/net_ib_cast/gin.cc @@ -418,6 +418,13 @@ ncclResult_t IbCastRmaIbProxyCreateContext(void* collComm, ncclRmaConfig_t* conf NCCLCHECKGOTO(ncclIbMalloc((void**)&gc->fullSendComm, sizeof(void*) * nranks), ret, end); NCCLCHECKGOTO(ncclIbMalloc((void**)&gc->fullRecvComm, sizeof(void*) * nranks), ret, end); gc->rank = cComm->rank; + // [RCCL] On AINIC a QP's channel id picks its UDMA engine set (connect.cc): the first QP of + // each (NIC, channel id) takes the set after the last one handed out, so new ids alternate. + // Without an id every context's QPs got channel 0 and all of GIN shared one set: ~30 GB/s per + // NIC, against 40-45 GB/s for RCCL send/recv over 8 channels. Ids from the top of the range, + // which the NET transport's channels never reach, make consecutive contexts alternate. + ncclNet_ctxt_t chCtxt = {}; + chCtxt.chId = (MAXCHANNELS - 1) - (c % (MAXCHANNELS / 2)); for (int i = 0; i < nranks; i += config->rankStride) { int connectPeer = (cComm->rank + i) % nranks; @@ -425,11 +432,12 @@ ncclResult_t IbCastRmaIbProxyCreateContext(void* collComm, ncclRmaConfig_t* conf do { if (gc->fullSendComm[connectPeer] == NULL) NCCLCHECKGOTO(IbCastConnectImpl(cComm->ctx, cComm->dev, handles + NCCL_NET_HANDLE_MAXSIZE * connectPeer, - &gc->fullSendComm[connectPeer], NULL, + &gc->fullSendComm[connectPeer], (ncclNetDeviceHandle_t**)&chCtxt, ncclParamCastGinIbTc() != -1 ? ncclParamCastGinIbTc() : ncclParamIbCastTc()), ret, end); if (gc->fullRecvComm[acceptPeer] == NULL) - NCCLCHECKGOTO(netIbCast.accept(lComm, &gc->fullRecvComm[acceptPeer], NULL), ret, end); + NCCLCHECKGOTO(netIbCast.accept(lComm, &gc->fullRecvComm[acceptPeer], (ncclNetDeviceHandle_t**)&chCtxt), + ret, end); } while ((gc->fullSendComm[connectPeer] == NULL) || (gc->fullRecvComm[acceptPeer] == NULL)); NCCLCHECKGOTO(IbCastGinIbP2PBarrier(cComm), ret, end); }