diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index b587fe24c270..bedc84fe748e 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -22,6 +22,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `NCCL_HIER_CE_COLL_AG_RAIL_RING_ENABLE` (default `-1`): a positive value selects a ring for the inter-node rail phase of hierarchical Copy Engine `ncclAllGather`. The default keeps the direct path. * `NCCL_IB_SORT_MERGE_NICS`: sorts the sub-devices of a merged IB device by plane ID. RCCL defaults it to `0` (NCCL defaults to `1`), so merged-device order and names are unchanged. * nccl4py: per-call collective configuration (`NCCLCollConfig`, `VendorOption`) including the launch completion event, communicator properties (`NCCLCommProperties`), and the `GIN_ONLY` window flag. On ROCm, the HIP `Event` shim provides only an event handle; `record()`, `sync()` and `query` are not implemented. +* `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged. ### Changed * Host-side device API declarations (`ncclDevCommCreate`, `ncclDevCommDestroy`, `ncclCommQueryProperties`, `ncclGetPeerDevicePointer`, the `*CreateRequirement` helpers, the `ncclDevCommRequirements` and `ncclCommProperties` structs and their initializers) moved to `nccl_device/host.h`. `nccl_device.h` still includes it; code that includes individual `nccl_device/*.h` headers directly must also include `nccl_device/host.h`. @@ -33,6 +34,13 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Network devices for each GPU are now chosen by rail and plane assignment, replacing the previous start-device scattering. * `NCCL_MLOPART_RDMA_ENABLE` (default `0`) is kept. NCCL 2.32 removed it and treats buffers on partitioned GPUs as RDMA-capable on all non-ARM hosts; RCCL keeps network buffer registration for partitioned (CPX/DPX) GPUs opt-in. * Retuned the symmetric AllReduce, ReduceScatter and AllGather kernels for gfx950. The block width is now selected per collective and message size instead of a fixed 256 threads, and the work partitioning and unroll factors in the load-direct kernels and the AllGather store kernel were refitted to the 64-lane wavefront. AllGather also switches from LL to the store kernel from 4 MB instead of 16 MB. The multi-node GIN kernels and all other architectures are unchanged. To fit the wider LL blocks, the shared LL scratch buffer on gfx950 grows from 4 MiB to 8 MiB per 8-rank communicator. +* gfx1250 no longer forces unregistered `ncclAllGather`, `ncclAlltoAll`, `ncclGather` and `ncclScatter` through the Copy Engine scratch path (CE-Scratch), because it was slower than the kernel paths it replaced. Whether CE-Scratch is forced by default now comes from the per-architecture table: gfx942 and gfx950 still force it. A communicator-level `NCCL_CTA_POLICY_ZERO` still forces it on every architecture; an explicit `RCCL_FORCE_CE=1` still forces it and `RCCL_FORCE_CE=0` still turns it off. gfx1250 `ncclAllGather` now uses DDA LL128 up to 128 MiB. +* gfx1250 `ncclAllReduce` on single-node communicators keeps the LL128 protocol at larger sizes: Tree at every size, Ring up to nRanks x 32 MiB, and the DDA LL128 tier up to 256 MiB. The DDA LL128 two-shot ceiling now follows that tier ceiling (`RCCL_DDA_LL128_TWOSHOT_THRESHOLD` still overrides it); the DDA scratch slot lets two-shot reach 256 MiB from 32 ranks, about 240 MiB at 16 and 128 MiB at 2 to 8. On gfx1250 communicators of 2 to 8 ranks, the DDA fabric scratch buffer grows to about 546 MiB. +* On gfx1250 single-node communicators of 8 or more ranks, the DDA tier, algorithm and protocol crossovers now scale with the rank count. Communicators below 8 ranks keep the previous thresholds. + * `ncclAllReduce`: DDA LL two-shot up to nRanks x 1 MiB (`RCCL_DDA_LL_TWOSHOT_THRESHOLD` still overrides it); from 32 ranks, DDA VMM between 2 MiB and nRanks x 512 KiB for unregistered buffers unless `RCCL_DDA_THRESHOLD` or `RCCL_DDA_LL_THRESHOLD` is set; Tree instead of Ring up to nRanks x 1 MiB unless `NCCL_ALGO` or `NCCL_PROTO` is set, with LL up to nRanks x 32 KiB. + * `ncclAllGather`: DDA LL up to nRanks x 16 KiB, DDA LL128 up to 256 MiB from 16 ranks, and Ring LL128 up to 64 MiB per rank. + * `ncclReduceScatter`: DDA LL up to nRanks x 256 KiB, DDA LL128 up to nRanks x 512 KiB, then DDA VMM up to nRanks x 2 MiB; Ring LL up to 512 KiB per rank and LL128 up to 8 MiB per rank (32 MiB for 1-byte types); Ring Simple on at most 128 channels between 8 and 32 MiB per rank unless `NCCL_MIN_NCHANNELS` or `NCCL_MAX_NCHANNELS` is set. + * DDA VMM `ncclAllReduce` and `ncclReduceScatter` launch 2048 / nRanks blocks, clamped to 32 to 128 and never more than before. ### Removed * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it. @@ -45,6 +53,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed tuner plugins receiving uninitialized cost-model constants. * Fixed GIN proxy descriptor shared-memory sizing and alignment. * Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly. +* Fixed Copy Engine `ncclAlltoAllv` hanging on intra-node communicators of more than 8 ranks with uneven send sizes. Each rank decided from its own send sizes whether to add intra-batch barriers, so ranks could issue different barrier counts. The decision now uses the full send-size matrix, which every rank shares. * Fixed `NCCL_P2P_DISABLE` and `NCCL_P2P_LEVEL` being ignored on Intel, ARM and Zhaoxin hosts, where the CPU-based default P2P level overwrote the user setting. On such hosts `NCCL_P2P_DISABLE=1` now falls back to SHM as documented, and `NCCL_P2P_LEVEL` is honored. ### Known issues diff --git a/projects/rccl/docs/api-reference/env-variables.rst b/projects/rccl/docs/api-reference/env-variables.rst index 1b730f2e265f..64c14bb1ceb0 100644 --- a/projects/rccl/docs/api-reference/env-variables.rst +++ b/projects/rccl/docs/api-reference/env-variables.rst @@ -825,6 +825,19 @@ threshold table introduced for gfx1250. | ``1``: Ignore the table; fall back to pre-table compile-time constants. This matches the behavior before the arch table was introduced. + * - | ``RCCL_TUNING_PARTITION_MODE`` + | Selects which gfx1250 per-architecture table (SPX or DPX) is used for + tuning thresholds. It only changes tuning, not results. A communicator + uses the DPX table only when every rank selects DPX; otherwise every rank + uses the SPX table. Other architectures have one table for every mode. + Has no effect when ``RCCL_IGNORE_ARCH_TABLE=1``. + - | ``auto`` (default): Follow the compute-partition mode the GPU reports. + ``DPX`` selects the DPX table; ``SPX``, ``CPX``, or no reported mode + selects the SPX table. Any other value is reported once and treated + as ``auto``. + | ``SPX``: Use the SPX table. + | ``DPX``: Use the DPX table. + * - | ``RCCL_DDA_THRESHOLD`` | Upper bound in bytes for the DDA VMM/Simple tier per collective. Messages above this size exit DDA and fall through to Ring, CE, or the diff --git a/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric.cu b/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric.cu index 25f4fc7102aa..8482499bce7d 100644 --- a/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric.cu +++ b/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric.cu @@ -13,6 +13,7 @@ #include "debug.h" #include "algorithms/dda/dda_init_detail.h" #include "algorithms/dda/fabric/fabric_gpu_barrier.h" +#include "rccl_common.h" #include #include @@ -29,9 +30,9 @@ using nccl_dda_detail::DdaFabricBarrierState; constexpr size_t kDdaFlatTreeThresholdBytes = 1ULL << 18; // Single source of the launch geometry: grid/block for `count` elements of -// `typeSize` bytes, capped at the init-time block count. +// `typeSize` bytes, capped at rcclDdaFabricVmmBlocks() (at most the init-time block count). static inline std::pair ddaAllReduceFabricGeom(ncclComm* comm, size_t count, int typeSize) { - return dda::common::getGridAndBlockDims(count, typeSize, comm->ddaFabricMaxBlocks); + return dda::common::getGridAndBlockDims(count, typeSize, rcclDdaFabricVmmBlocks(comm)); } template @@ -57,7 +58,7 @@ static ncclResult_t ncclAllReduceDdaFabricTyped(const void* sendbuff, void* recv count, nRanks); } - // Block cap chosen at init (barrier flag buffer is sized for it). + // At most the block cap chosen at init (barrier flag buffer is sized for it). auto gridBlock = ddaAllReduceFabricGeom(comm, count, sizeof(T)); const auto& grid = gridBlock.first; const auto& block = gridBlock.second; diff --git a/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric_ll.cu b/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric_ll.cu index f6870fa12937..cc67a7e83d89 100644 --- a/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric_ll.cu +++ b/projects/rccl/src/algorithms/dda/all_reduce/dda_all_reduce_fabric_ll.cu @@ -26,6 +26,7 @@ #include "debug.h" #include "algorithms/dda/fabric/fabric_gpu_barrier.h" // dda::common::kDdaMaxNranks #include "param.h" +#include "rccl_common.h" #include #include @@ -37,9 +38,7 @@ RCCL_PARAM_DECLARE(DdaLL); RCCL_PARAM_DECLARE(DdaLLOneShotThreshold); -RCCL_PARAM_DECLARE(DdaLLTwoShotThreshold); RCCL_PARAM_DECLARE(DdaLL128OneShotThreshold); -RCCL_PARAM_DECLARE(DdaLL128TwoShotThreshold); namespace { @@ -366,7 +365,7 @@ bool ddaLLArTwoShotEligible(ncclComm* comm, const void* sendbuff, void* recvbuff return false; } - if (count * ncclTypeSize(datatype) > (size_t)rcclParamDdaLLTwoShotThreshold()) { + if (count * ncclTypeSize(datatype) > rcclDdaLLTwoShotThreshold(comm)) { return false; } @@ -471,7 +470,7 @@ bool ddaLL128ArTwoShotEligible(ncclComm* comm, const void* sendbuff, void* recvb return false; } - if (count * ncclTypeSize(datatype) > (size_t)rcclParamDdaLL128TwoShotThreshold()) { + if (count * ncclTypeSize(datatype) > rcclDdaLL128TwoShotThreshold(comm)) { return false; } diff --git a/projects/rccl/src/algorithms/dda/reduce_scatter/dda_reduce_scatter_fabric.cu b/projects/rccl/src/algorithms/dda/reduce_scatter/dda_reduce_scatter_fabric.cu index 4724c2de6fe9..b8b06ab96f0e 100644 --- a/projects/rccl/src/algorithms/dda/reduce_scatter/dda_reduce_scatter_fabric.cu +++ b/projects/rccl/src/algorithms/dda/reduce_scatter/dda_reduce_scatter_fabric.cu @@ -13,6 +13,7 @@ #include "debug.h" #include "algorithms/dda/dda_init_detail.h" #include "algorithms/dda/fabric/fabric_gpu_barrier.h" +#include "rccl_common.h" #include #include @@ -40,8 +41,8 @@ static ncclResult_t ncclReduceScatterDdaFabricTyped(const void* sendbuff, void* return ncclInvalidArgument; } - // Use the block cap chosen at init (barrier flag buffer is sized for it). - const int nBlocksMax = comm->ddaFabricMaxBlocks; + // At most the block cap chosen at init (barrier flag buffer is sized for it). + const int nBlocksMax = rcclDdaFabricVmmBlocks(comm); // For reduce-scatter we use recvcount for grid calculation since each rank // processes its own shard. auto gridBlock = dda::common::getGridAndBlockDims(recvcount, sizeof(T), nBlocksMax); diff --git a/projects/rccl/src/ce_coll.cc b/projects/rccl/src/ce_coll.cc index e7f011a78c7a..27a168127b79 100644 --- a/projects/rccl/src/ce_coll.cc +++ b/projects/rccl/src/ce_coll.cc @@ -1150,17 +1150,38 @@ ncclResult_t ncclCeAlltoAllv(struct ncclComm* comm, struct ncclCeCollArgs* args, uint8_t* myRecvBuff = (uint8_t*)args->recvBuff; void* peerRecvBuff; size_t offset, winOff; - size_t totalBytes = 0; + const int freq = (int)comm->ceColl.intraBatchSyncFreq; + bool roundSync = false; struct ncclCeBatchOpsParams batchOpsParams = {}; NCCLCHECKGOTO(ncclCeInitBatchOpsParams(&batchOpsParams, comm->nRanks), ret, fail); NCCLCHECKGOTO(ncclAlltoAllvValidateSizeMatrix(args->sizes, comm->nRanks), ret, fail); + // ncclMemOpSync is collective: decide intra-batch sync from the full size matrix, not this rank's nonzero peers. + for (int src = 0; freq > 0 && src < comm->nRanks && !roundSync; src++) { + size_t* srcSizes = ncclAlltoAllvSendSizes(args->sizes, src, comm->nRanks); + int ops = 0; + size_t bytes = 0; + for (int dst = 0; dst < comm->nRanks; dst++) { + if (srcSizes[dst] != 0) { + ops++; + bytes += srcSizes[dst]; + } + } + roundSync = ops > freq && bytes >= comm->ceColl.intraBatchSyncMsgThreshold; + } + // Ensure all ranks are ready before starting transfers NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail); // Copy data to other ranks: send variable-sized chunk for each destination rank for (int r = 0; r < comm->nRanks; r++) { + // Sync per round of freq destination slots, so every rank issues the same number of barriers. + if (roundSync && r > 0 && r % freq == 0) { + NCCLCHECKGOTO(ncclCeLaunchBatchOps(comm, &batchOpsParams, stream, args), ret, fail); + batchOpsParams.numOps = 0; + NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail); + } int dstRank = (comm->rank + r) % comm->nRanks; const size_t chunkBytes = sendSizes[dstRank]; if (chunkBytes == 0) { @@ -1169,7 +1190,6 @@ ncclResult_t ncclCeAlltoAllv(struct ncclComm* comm, struct ncclCeCollArgs* args, uint8_t* srcPtr = mySendBuff + sendDispls[dstRank]; uint8_t* dstPtr = myRecvBuff + recvDispls[comm->rank]; - totalBytes += chunkBytes; if (dstRank == comm->rank) { // Local copy for own data @@ -1197,13 +1217,14 @@ ncclResult_t ncclCeAlltoAllv(struct ncclComm* comm, struct ncclCeCollArgs* args, } } - // Check if we need to perform intra-batch synchronization - batchOpsParams.intraBatchSync = - (batchOpsParams.numOps > comm->ceColl.intraBatchSyncFreq && totalBytes >= comm->ceColl.intraBatchSyncMsgThreshold); - - // Launch the batch operations + // Launch the last round; intraBatchSync stays false so ncclCeLaunchBatchOps adds no rank-local barriers NCCLCHECKGOTO(ncclCeLaunchBatchOps(comm, &batchOpsParams, stream, args), ret, fail); + // Graph replay needs an even barrier count so useCompletePtr returns to its pre-launch state. + if (roundSync && ncclCudaGraphValid(comm->planner.capturingGraph) && ((comm->nRanks - 1) / freq) % 2 != 0) { + NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail); + } + // Ensure all transfers are complete across all ranks NCCLCHECKGOTO(ncclMemOpSync(comm, stream, args), ret, fail); diff --git a/projects/rccl/src/enqueue/enqueue.cc b/projects/rccl/src/enqueue/enqueue.cc index 750a98d2f50b..2df31477d30f 100644 --- a/projects/rccl/src/enqueue/enqueue.cc +++ b/projects/rccl/src/enqueue/enqueue.cc @@ -3070,6 +3070,7 @@ static ncclResult_t topoGetAlgoInfo(struct ncclComm* comm, struct ncclTaskColl* } // Honor Tuner config if available if (!isTunerMatchFound && info->algorithm != NCCL_ALGO_PAT) { + rcclUpdateCollectiveAlgorithm(comm, nBytes, info, table); rcclUpdateCollectiveProtocol(comm, nBytes, info); } rcclSetPipelining(comm, nBytes, info); @@ -3122,6 +3123,7 @@ static ncclResult_t topoGetAlgoInfo(struct ncclComm* comm, struct ncclTaskColl* } INFO(NCCL_TUNING, "post-adjustment based on threadThreshold:%i nBytes:%lu nc:%i", threadThreshold, nBytes, nc); rcclOverrideChannels(comm, info->func, nBytes, nc); + rcclUpdateCollectiveChannels(comm, nBytes, info, nc); } #ifdef ENABLE_ROCSHMEM @@ -4765,8 +4767,8 @@ static ncclResult_t taskAppend(struct ncclComm* comm, struct ncclInfo* info) { opDev)); } else if (!rcclNcclAlgoEnvIsSet() && !allGatherDecided && !alltoAllDecided && - rcclParamForceCe() && CeScratchAvailable && !hasSysmemSegment && comm->ddaScratch != nullptr && - recvBytes <= comm->ddaScratchBytes && info->coll != ncclFuncAllReduce) { + rcclForceCeScratchEnabled(comm) && CeScratchAvailable && !hasSysmemSegment && + comm->ddaScratch != nullptr && recvBytes <= comm->ddaScratchBytes && info->coll != ncclFuncAllReduce) { INFO(NCCL_TUNING, "Using DDA scratch for CE collective, count=%zu, recvBytes=%zu", info->count, recvBytes); NCCLCHECK(ceCollTaskAppend(comm, info, /*sendWin=*/nullptr, /*recvWin=*/nullptr, comm->ddaScratch, comm->ddaPeerPtrsHost, opDev)); diff --git a/projects/rccl/src/include/comm.h b/projects/rccl/src/include/comm.h index 5861387ead31..22ddcd882b93 100644 --- a/projects/rccl/src/include/comm.h +++ b/projects/rccl/src/include/comm.h @@ -645,6 +645,9 @@ struct ncclPeerInfo { bool gpuCftMulticastSupport; bool gpuCftCountedSupport; uint32_t gitVersionHash; + // [RCCL] The rcclTuningPartition this rank asks for (RCCL_TUNING_PARTITION_MODE). The comm + // uses DPX only when every rank asks for it. Append later fields after this one. + int tuningPartition; }; typedef enum ncclGroupTaskType { diff --git a/projects/rccl/src/include/rccl_arch_thresholds.h b/projects/rccl/src/include/rccl_arch_thresholds.h index 5093e0e95fdd..86aae7927d06 100644 --- a/projects/rccl/src/include/rccl_arch_thresholds.h +++ b/projects/rccl/src/include/rccl_arch_thresholds.h @@ -79,8 +79,64 @@ struct rcclArchThresholds { // 0 means no suppression (symk may win at any size for that collective). size_t symMinR2[RCCL_DDA_FUNC_COUNT]; + // Rank-scaled tier caps, bytes per rank. On a comm of at least + // rankScaledMinRanks ranks, a non-zero entry replaces the matching + // ddaLLMax / ddaLL128Max / ddaVmmMax total with nRanks * entry, clipped to + // rankScaledMaxBytes (which must not exceed the largest total above: fabric + // scratch is sized from the totals). Smaller comms, and entries of 0, keep the + // totals. The RCCL_DDA_*_THRESHOLD env vars still win over both. + size_t ddaLLMaxPerRank[RCCL_DDA_FUNC_COUNT]; + size_t ddaLL128MaxPerRank[RCCL_DDA_FUNC_COUNT]; + size_t ddaVmmMaxPerRank[RCCL_DDA_FUNC_COUNT]; + int rankScaledMinRanks; // 0 = no rank scaling + size_t rankScaledMaxBytes; + + // AllReduce only: on comms of at least ddaArVmmFirstMinRanks ranks, DDA VMM + // takes unregistered messages in (ddaArVmmFirstMinBytes, nRanks * + // ddaArVmmFirstMaxPerRank] ahead of the LL tier, provided the VMM path is + // eligible. 0 ranks = off. + int ddaArVmmFirstMinRanks; + size_t ddaArVmmFirstMinBytes; + size_t ddaArVmmFirstMaxPerRank; + + // DDA VMM (AllReduce / ReduceScatter) launches min(init block cap, + // clamp(ddaVmmBlockBudget / nRanks, ddaVmmMinBlocks, ddaVmmMaxBlocks)) blocks + // on comms of at least rankScaledMinRanks ranks. Budget 0 = the init block cap. + int ddaVmmBlockBudget; + int ddaVmmMinBlocks; + int ddaVmmMaxBlocks; + + // Ring / Tree crossovers that replace the tuning model's ranges on single-node + // comms (nNodes == 1). Sizes are per rank as rcclGetSizePerRank() reports them + // (AllReduce: the whole message, hence nRanks x the entry). 0 = keep the + // model's range. The two AllReduce LL128 entries apply at any rank count, the + // rest from rankScaledMinRanks ranks. + size_t arRingLL128MaxPerRank; // Ring AllReduce LL128 up to nRanks x this + bool arTreeLL128Unlimited; // Tree AllReduce keeps LL128 at every size + size_t arTreeMaxPerRank; // AllReduce runs Tree instead of Ring up to nRanks x this, + size_t arTreeLLMaxPerRank; // with LL up to nRanks x this and LL128 above + size_t agRingLL128MaxPerRank; // AllGather Ring LL128 up to this per rank + size_t rsRingLL128MinPerRank; // ReduceScatter Ring LL128 above this per rank, + size_t rsRingLL128MaxPerRank; // up to this (2- and 4-byte types) + size_t rsRingLL128MaxPerRank1B; // or this (1-byte types) + size_t rsRingSimpleCappedMaxPerRank; // Ring Simple ReduceScatter above the LL128 band, + int rsRingSimpleCappedChannels; // up to this per rank, runs on at most this many channels + + // Force CE-Scratch for unregistered AllGather / AlltoAll / Gather / Scatter + // whenever the receive fits ddaScratch, on top of the ceNonRegMin/Max window. + // NCCL_CTA_POLICY_ZERO forces it on every arch (it is the one zero-CU path for + // unregistered buffers). RCCL_FORCE_CE=0 turns it off everywhere. A null table forces it. + bool forceCeScratch; +}; + +// Compute-partition mode a gfx1250 comm is tuned for (RCCL_TUNING_PARTITION_MODE). It picks +// the table only; other archs have one table for every mode. +enum rcclTuningPartition { + RCCL_TUNING_PARTITION_SPX = 0, + RCCL_TUNING_PARTITION_DPX = 1, }; -const rcclArchThresholds* rcclGetArchThresholds(const char* gcn); +const rcclArchThresholds* rcclGetArchThresholds(const char* gcn, + rcclTuningPartition partition = RCCL_TUNING_PARTITION_SPX); #endif // RCCL_ARCH_THRESHOLDS_H_ diff --git a/projects/rccl/src/include/rccl_common.h b/projects/rccl/src/include/rccl_common.h index 12fc4de0ba2f..062c3370bd2e 100644 --- a/projects/rccl/src/include/rccl_common.h +++ b/projects/rccl/src/include/rccl_common.h @@ -146,6 +146,15 @@ ncclResult_t rcclOverrideProtocol(const char* ncclProtoStr[], float table[][NCCL ncclResult_t rcclOverrideAlgorithm(const char* ncclAlgoStr[], float table[][NCCL_NUM_PROTOCOLS], struct ncclTaskColl* info); void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info); +// Arch-table AllReduce Tree-over-Ring switch (rcclArchThresholds::arTreeMaxPerRank; +// single-node comms of rankScaledMinRanks+ ranks, NCCL_ALGO / NCCL_PROTO unset, cost +// table offering Tree LL and LL128). Call before rcclUpdateCollectiveProtocol, which +// then picks the Tree protocol. +void rcclUpdateCollectiveAlgorithm(struct ncclComm* comm, size_t nBytes, struct ncclTaskColl* info, + float table[][NCCL_NUM_PROTOCOLS]); +// Arch-table channel cap for Ring Simple ReduceScatter +// (rcclArchThresholds::rsRingSimpleCapped*). NCCL_MIN/MAX_NCHANNELS turn it off. +void rcclUpdateCollectiveChannels(struct ncclComm* comm, size_t nBytes, const struct ncclTaskColl* info, int& nc); void rcclUpdateThreadThreshold(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info, int& threadThreshold); void rcclSetPipelining(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info); @@ -253,6 +262,10 @@ bool rcclCeAllReduceEnabled(const struct ncclComm* comm); // Same resolution for RCCL_FORCE_CE_ALLREDUCE, which lets CE AllReduce (2-shot // and registered) run without CTAPolicy=ZERO. bool rcclForceCeAllReduceEnabled(const struct ncclComm* comm); +// Is CE-Scratch forced for unregistered AllGather / AlltoAll / Gather / Scatter? +// RCCL_FORCE_CE=0 turns it off; otherwise NCCL_CTA_POLICY_ZERO forces it and +// the arch table's forceCeScratch decides. +bool rcclForceCeScratchEnabled(const struct ncclComm* comm); RCCL_PARAM_DECLARE(CeArMaxMsgBytes); // -1 = use ceNonRegMax[AR] (2-shot) from arch table RCCL_PARAM_DECLARE(CeArRegMaxMsgBytes); // -1 = use ceRegMax[AR] (registered) from arch table RCCL_PARAM_DECLARE(CeArStagingBytes); // -1 = use NCCL_CE_AR_STAGING_BYTES; sizes ceARTmpBuf, not the cap @@ -330,6 +343,16 @@ constexpr size_t kDdaVmmBaseDefault = 134217728; // 128 MiB size_t rcclDdaLLThreshold(const ncclComm* comm, ncclFunc_t func); size_t rcclDdaLL128Threshold(const ncclComm* comm, ncclFunc_t func); size_t rcclDdaVmmThreshold(const ncclComm* comm, ncclFunc_t func); +// Ceiling for the DDA LL128 two-shot AllReduce: RCCL_DDA_LL128_TWOSHOT_THRESHOLD +// when set, else the AllReduce LL128 tier cap where the arch table has that tier, +// else the env var's default. +size_t rcclDdaLL128TwoShotThreshold(const ncclComm* comm); +// Ceiling for the DDA LL two-shot AllReduce: RCCL_DDA_LL_TWOSHOT_THRESHOLD when +// set; on rank-scaled comms the AllReduce LL tier cap; else the env var's default. +size_t rcclDdaLLTwoShotThreshold(const ncclComm* comm); +// Blocks the DDA VMM AllReduce / ReduceScatter kernels launch: the init cap, +// lowered on large comms per rcclArchThresholds::ddaVmmBlockBudget. +int rcclDdaFabricVmmBlocks(const ncclComm* comm); // Cap for the DDA entry gate: the max of the LL, LL128 and VMM caps above, // skipping tiers turned off by RCCL_DDA_LL / RCCL_DDA_LL128. Pass this to diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 092f6a2fe3ee..d2a8fd105c98 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -1264,6 +1264,46 @@ static const char* mnnvlDegradedBwStr(unsigned int healthMask) { } #endif +// RCCL_TUNING_PARTITION_MODE: the gfx1250 tuning table this rank asks for. "auto" (or unset) +// follows the compute-partition mode the GPU reports; "SPX" / "DPX" force one. It only picks +// tuning thresholds. Any mode other than DPX (CPX, or none reported) asks for SPX. An invalid +// value is reported once and treated as "auto". +static rcclTuningPartition rcclTuningPartitionRequest(const char* computePartition) { + const char* mode = ncclGetEnv("RCCL_TUNING_PARTITION_MODE"); + if (mode != nullptr && mode[0] != '\0' && strcasecmp(mode, "auto") != 0) { + if (strcasecmp(mode, "SPX") == 0) return RCCL_TUNING_PARTITION_SPX; + if (strcasecmp(mode, "DPX") == 0) return RCCL_TUNING_PARTITION_DPX; + static std::once_flag invalidModeOnce; + std::call_once(invalidModeOnce, [mode] { + ATTN("Invalid value %s for RCCL_TUNING_PARTITION_MODE (expected auto, SPX or DPX), using auto", mode); + }); + } + return computePartition != nullptr && strcasecmp(computePartition, "DPX") == 0 ? RCCL_TUNING_PARTITION_DPX + : RCCL_TUNING_PARTITION_SPX; +} + +// Ranks pick algorithms on their own, so every rank of a comm must read the same table: DPX only +// when all of them ask for it, SPX otherwise. *mixed reports ranks that disagree. +static rcclTuningPartition rcclTuningPartitionForComm(const struct ncclPeerInfo* peers, int nRanks, bool* mixed) { + int nDpx = 0; + for (int i = 0; i < nRanks; i++) nDpx += peers[i].tuningPartition == RCCL_TUNING_PARTITION_DPX; + *mixed = nDpx > 0 && nDpx < nRanks; + return nDpx == nRanks ? RCCL_TUNING_PARTITION_DPX : RCCL_TUNING_PARTITION_SPX; +} + +// Re-picks comm->archThresholds once every rank's request is known (after AllGather1). +static void rcclSetTuningPartition(struct ncclComm* comm) { + bool mixed = false; + const rcclTuningPartition partition = rcclTuningPartitionForComm(comm->peerInfo, comm->nRanks, &mixed); + comm->archThresholds = rcclGetArchThresholds(comm->archName, partition); + if (comm->rank != 0 || !IsArchMatch(comm->archName, "gfx1250")) return; + if (mixed) { + ATTN("gfx1250 tuning table: ranks ask for different tables (RCCL_TUNING_PARTITION_MODE or GPU partition mode), " + "using SPX on every rank"); + } + INFO(NCCL_INIT | NCCL_TUNING, "gfx1250 tuning table: %s", partition == RCCL_TUNING_PARTITION_DPX ? "DPX" : "SPX"); +} + static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, uint64_t commHash) { cudaDeviceProp prop; info->rank = comm->rank; @@ -1311,18 +1351,19 @@ static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, u // partition 0 are both function 0 with an accelerator class, so read the mode from sysfs and // leave mloPart undefined on an unpartitioned GPU, which must not get the DEV overlay (it breaks // Rome gpuId matching and disables GIN/GDR). + // The partition mode lives on the physical device. A CPX alias at .1-.7 is usually absent + // from sysfs entirely, so ask function 0 rather than our own BDF. + char physBusIdStr[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE]; + char partition[MAX_STR_LEN]; + partition[0] = '\0'; + if (int64ToBusId(info->busId & ~0xfLL, physBusIdStr) == ncclSuccess) { + (void)ncclOsGetPciDeviceComputePartitionByBusId(physBusIdStr, partition, sizeof(partition)); + } + info->tuningPartition = rcclTuningPartitionRequest(partition); if (info->mloPart == NCCL_TOPO_UNDEF) { int fn = (int)(info->busId & 0xf); if (fn < NCCL_TOPO_MLOPART_DEV_MAX) { char busIdStr[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE]; - // The partition mode lives on the physical device. A CPX alias at .1-.7 is usually absent - // from sysfs entirely, so ask function 0 rather than our own BDF. - char physBusIdStr[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE]; - char partition[MAX_STR_LEN]; - partition[0] = '\0'; - if (int64ToBusId(info->busId & ~0xfLL, physBusIdStr) == ncclSuccess) { - (void)ncclOsGetPciDeviceComputePartitionByBusId(physBusIdStr, partition, sizeof(partition)); - } // A partitioned device makes every function a partition, function 0 included: partition 0 is // a partition, not an unpartitioned GPU, and must carry index 0 so all of a device's // partitions share one DEV overlay group. An empty string means the platform does not report @@ -1854,6 +1895,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p } } } + rcclSetTuningPartition(comm); // AllGather1 - end timers[TIMER_INIT_ALLGATHER] = clockNano() - timers[TIMER_INIT_ALLGATHER]; @@ -3111,6 +3153,7 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) { comm->ll128LineElems = rcclLL128LineElemsFromArch(comm->archName); comm->ll128DataElems = rcclLL128DataElemsFromArch(comm->archName); comm->ll128ShmemElemsPerThread = rcclLL128ShmemElemsPerThreadFromArch(comm->archName); + // SPX until initTransportsRank knows every rank's RCCL_TUNING_PARTITION_MODE request (rcclSetTuningPartition). comm->archThresholds = rcclGetArchThresholds(comm->archName); comm->maxSharedMemOptin = maxSharedMem; diff --git a/projects/rccl/src/rccl_arch_thresholds.cc b/projects/rccl/src/rccl_arch_thresholds.cc index 6e9616708518..13b961ee8c98 100644 --- a/projects/rccl/src/rccl_arch_thresholds.cc +++ b/projects/rccl/src/rccl_arch_thresholds.cc @@ -21,6 +21,9 @@ // All collectives compare total message bytes against the table values. // gfx1250 placeholders -- validate against sweep data (AICOMRCCL-1756). +// SPX table, also used for CPX and when no partition mode is reported. DPX has its own +// copy below (rcclArchThresholds_gfx1250_dpx): a non-tuning change here (a gate such as +// forceCeScratch, a disabled slot) usually belongs in both. // Index mapping: [Bcast=0, Reduce=1, AG=2, RS=3, AR=4, SR=5, Send=6, Recv=7, A2A=8] static const rcclArchThresholds rcclArchThresholds_gfx1250 = { // ddaLLMax: DDA LL tier ceiling per collective (fabric gfx1250 only). @@ -40,9 +43,10 @@ static const rcclArchThresholds rcclArchThresholds_gfx1250 = { .ddaLL128Max = { 0, // [0] Broadcast -- not used 0, // [1] Reduce -- not used - 64ULL*1024*1024, // [2] AllGather + 128ULL*1024*1024, // [2] AllGather -- beats Ring at 128M on 8-64 ranks (AICOMRCCL-2876) 2ULL*1024*1024, // [3] ReduceScatter - 32ULL*1024*1024, // [4] AllReduce + 256ULL*1024*1024, // [4] AllReduce -- two-shot beats Ring to 256M (AICOMRCCL-2875); its scratch + // slot reaches 256M from 32 ranks, ~128M at 2-8, 240M at 16 0, // [5] SendRecv -- not used 0, // [6] Send -- not used 0, // [7] Recv -- not used @@ -170,6 +174,294 @@ static const rcclArchThresholds rcclArchThresholds_gfx1250 = { kThreshUnlimited, // [7] Recv -- not used on gfx1250 kThreshUnlimited, // [8] AlltoAll -- not used on gfx1250 }, + // Rank-scaled caps from the 8-64 rank H13 sweep, AICOMRCCL-1756. The + // crossovers move with the per-rank share, so on 8+ rank comms these replace the + // totals above (which stay in force for smaller comms). + .ddaLLMaxPerRank = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 16ULL*1024, // [2] AllGather -- LL beats LL128 to 128K/256K/512K/1M at 8/16/32/64 + 256ULL*1024, // [3] ReduceScatter -- LL past 2M/4M/8M/16M at 8/16/32/64 trails LL128 up to 2x + 1ULL*1024*1024, // [4] AllReduce -- LL two-shot beats LL128 1.4-1.9x to 8M/16M/32M/64M + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- unchanged + }, + .ddaLL128MaxPerRank = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 16ULL*1024*1024, // [2] AllGather -- 128M at 8 ranks, 256M (clipped) from 16 + 512ULL*1024, // [3] ReduceScatter -- the LL128 slot reaches nRanks x 512K + 0, // [4] AllReduce -- ddaLL128Max total + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- unchanged + }, + .ddaVmmMaxPerRank = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- VMM off + 2ULL*1024*1024, // [3] ReduceScatter -- VMM beats Ring 1.2-1.9x to 16M/32M/64M/128M + 0, // [4] AllReduce -- ddaVmmMax total (see ddaArVmmFirst*) + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- VMM off + }, + .rankScaledMinRanks = 8, + .rankScaledMaxBytes = 256ULL*1024*1024, + // From 32 ranks, AllReduce VMM beats LL two-shot / LL128 1.3-3.7x between 2M and nRanks x 512K. + .ddaArVmmFirstMinRanks = 32, + .ddaArVmmFirstMinBytes = 2ULL*1024*1024, + .ddaArVmmFirstMaxPerRank = 512ULL*1024, + // VMM peaks at 128 / 128 / 64 / 32 blocks on 8 / 16 / 32 / 64 ranks. + .ddaVmmBlockBudget = 2048, + .ddaVmmMinBlocks = 32, + .ddaVmmMaxBlocks = 128, + // Ring / Tree crossovers, same sweep (AllReduce LL128: AICOMRCCL-2874/2875). + .arRingLL128MaxPerRank = 32ULL*1024*1024, // Ring LL128 beats Simple to 256M/512M/1G/2G at 8-64 + .arTreeLL128Unlimited = true, // Tree LL128 beats Tree Simple 1.3-2.9x at 32M-16G + .arTreeMaxPerRank = 1ULL*1024*1024, // Tree beats Ring 1.4-2.3x to 8M/16M/32M/64M + .arTreeLLMaxPerRank = 32ULL*1024, // Tree LL beats LL128 1.2-1.6x to 256K/512K/1M/2M + .agRingLL128MaxPerRank = 64ULL*1024*1024, // Ring LL128 beats Simple to 512M/1G/2G/4G + .rsRingLL128MinPerRank = 512ULL*1024, // Ring LL128 beats LL from 1M per rank + .rsRingLL128MaxPerRank = 8ULL*1024*1024, // 2/4-byte types: Simple wins past 8M per rank + .rsRingLL128MaxPerRank1B = 32ULL*1024*1024, // fp8: LL128 holds to 32M per rank + .rsRingSimpleCappedMaxPerRank = 32ULL*1024*1024,// 128 channels beat 192 by 1.1-1.3x in (8M, 32M] + .rsRingSimpleCappedChannels = 128, + // CE-Scratch trails the DDA / Ring / Direct kernels it displaces between the DDA + // ceiling and the scratch size (nRanks x 64 MiB) (AICOMRCCL-2876). + .forceCeScratch = false, +}; + +// gfx1250 in DPX mode (each GPU split into two partitions), picked by +// RCCL_TUNING_PARTITION_MODE. Started as a copy of the SPX table above; change entries +// here, not there, to tune DPX. Entries marked (DPX) were re-measured on DPX +// nodes at 8/16/32 ranks (rccl-tests -n 10, median of 3 rounds); the other per-entry +// notes are still the SPX measurements. +// Index mapping: [Bcast=0, Reduce=1, AG=2, RS=3, AR=4, SR=5, Send=6, Recv=7, A2A=8] +static const rcclArchThresholds rcclArchThresholds_gfx1250_dpx = { + // ddaLLMax: DDA LL tier ceiling per collective (fabric gfx1250 only). + .ddaLLMax = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 512ULL*1024, // [2] AllGather -- LL beats LL128 1.4-1.6x to 512K at 8 ranks; the + // crossover falls with rank count (512K at 16, ~128K + // at 32), so this is a total, not rank-scaled (DPX) + 16ULL*1024*1024, // [3] ReduceScatter + 32ULL*1024*1024, // [4] AllReduce + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 2ULL*1024*1024, // [8] AlltoAll -- LL beats Direct 1.2-4.3x to 2M at 8/16/32 (DPX) + }, + // ddaLL128Max: DDA LL128 tier ceiling per collective. + // RCCL_PARAM(DdaLL128, ...) defaults to 1 (enabled); set RCCL_DDA_LL128=0 to disable. + .ddaLL128Max = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 64ULL*1024*1024, // [2] AllGather -- Ring beats it 1.1-1.4x at 128M-256M on 8/16/32 (DPX) + 2ULL*1024*1024, // [3] ReduceScatter + 256ULL*1024*1024, // [4] AllReduce -- two-shot beats Ring to 256M (AICOMRCCL-2875); its scratch + // slot reaches 256M from 32 ranks, ~128M at 2-8, 240M at 16 + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 1ULL*1024*1024, // [8] AlltoAll + }, + // ddaVmmMax: DDA VMM (fabric simple) tier ceiling per collective. + // Messages above this fall to Ring/CE (or sym kernel for R2). + .ddaVmmMax = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- 0 = disabled + 16ULL*1024*1024, // [3] ReduceScatter -- wins only at 16M on 8 ranks; the LL tier covers it + // below 8 ranks and from 16 (DPX) + 16ULL*1024*1024, // [4] AllReduce + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- 0 = disabled + }, + // ddaVmmMaxR2: DDA VMM cap when recv buffer is registered (R2 mode). + // DDA is gated on !symEligible (AR: !symkRequested) on every arch, including + // gfx1250 fabric; R2 sum/avg never takes DDA. A non-zero entry only shortens + // DDA when recv is registered and that gate still passes. 0 means use ddaVmmMax. + .ddaVmmMaxR2 = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- not used + 0, // [3] ReduceScatter -- not used + 4ULL*1024*1024, // [4] AllReduce + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- not used + }, + // ddaVmmMaxGraph: DDA VMM cap during graph capture (graphCapturingHint=true). + // CE AllReduce is blocked by graphModeSeen latch during graph captures. + // Extending DDA VMM for AR lets DDA fill the window CE would otherwise absorb. + // 0 means use ddaVmmMax (no graph-specific override for that collective). + .ddaVmmMaxGraph = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- not used + 0, // [3] ReduceScatter -- not used + 32ULL*1024*1024, // [4] AllReduce + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- not used + }, + // ceNonRegMin: lower bound for CE-Scratch window per collective (-R 0). + // 0 = disabled for that collective. + .ceNonRegMin = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- 0 = disabled + 0, // [3] ReduceScatter -- not used + 0, // [4] AllReduce -- not used + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- not used + }, + // ceNonRegMax: upper bound for CE-Scratch window per collective (-R 0). + // AR: 0 = 2-shot off. ceARTmpBuf stays at the default 256 MiB unless this entry is larger. + .ceNonRegMax = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- 0 = disabled + 0, // [3] ReduceScatter -- not used + 0, // [4] AllReduce -- 0 = 2-shot off + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- not used + }, + // ceRegMax: registered CE upper bound per collective (-R 2). + .ceRegMax = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 8ULL*1024*1024*1024, // [2] AllGather + 0, // [3] ReduceScatter -- not used + 256ULL*1024*1024, // [4] AllReduce + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- not used + }, + // symMaxR2: suppress symk in favour of CE-registered when recv is registered and + // msg > threshold. kThreshUnlimited = no suppression; non-zero literal = byte cap; + // 0 = always suppress (symk not used for this collective on gfx1250). + .symMaxR2 = { + 0, // [0] Broadcast -- not used on gfx1250 + 0, // [1] Reduce -- not used on gfx1250 + 2ULL*1024*1024, // [2] AllGather + kThreshUnlimited, // [3] ReduceScatter + 0, // [4] AllReduce + 0, // [5] SendRecv -- not used on gfx1250 + 0, // [6] Send -- not used on gfx1250 + 0, // [7] Recv -- not used on gfx1250 + 0, // [8] AlltoAll -- not used on gfx1250 + }, + // Graph capture: CE is blocked, so symk or DDA wins. kThreshUnlimited = no suppression; + // non-zero literal = byte cap above which symk is withdrawn so DDA can win instead; + // 0 = always suppress (symk not used for this collective on gfx1250). + .symMaxR2Graph = { + 0, // [0] Broadcast -- not used on gfx1250 + 0, // [1] Reduce -- not used on gfx1250 + 2ULL*1024*1024, // [2] AllGather + kThreshUnlimited, // [3] ReduceScatter + 0, // [4] AllReduce + 0, // [5] SendRecv -- not used on gfx1250 + 0, // [6] Send -- not used on gfx1250 + 0, // [7] Recv -- not used on gfx1250 + 0, // [8] AlltoAll -- not used on gfx1250 + }, + // symMinR2: suppress symk below this size for R2 buffers so DDA wins in that sub-range. + // kThreshUnlimited = always suppress (totalBytes < SIZE_MAX always true) for unused collectives. + .symMinR2 = { + kThreshUnlimited, // [0] Broadcast -- not used on gfx1250 + kThreshUnlimited, // [1] Reduce -- not used on gfx1250 + 128ULL*1024, // [2] AllGather + 2ULL*1024*1024, // [3] ReduceScatter + 8ULL*1024*1024, // [4] AllReduce + kThreshUnlimited, // [5] SendRecv -- not used on gfx1250 + kThreshUnlimited, // [6] Send -- not used on gfx1250 + kThreshUnlimited, // [7] Recv -- not used on gfx1250 + kThreshUnlimited, // [8] AlltoAll -- not used on gfx1250 + }, + // Rank-scaled caps from the 8-64 rank H13 sweep, AICOMRCCL-1756. The + // crossovers move with the per-rank share, so on 8+ rank comms these replace the + // totals above (which stay in force for smaller comms). + .ddaLLMaxPerRank = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- ddaLLMax total (DPX) + 1ULL*1024*1024, // [3] ReduceScatter -- LL beats DDA Simple / Ring 1.1-1.5x to 8M/16M/32M + // at 8/16/32 (DPX) + 1ULL*1024*1024, // [4] AllReduce -- LL two-shot beats LL128 1.4-1.9x to 8M/16M/32M/64M + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- unchanged + }, + .ddaLL128MaxPerRank = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- ddaLL128Max total (DPX) + 512ULL*1024, // [3] ReduceScatter -- the LL128 slot reaches nRanks x 512K + 0, // [4] AllReduce -- ddaLL128Max total + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- unchanged + }, + .ddaVmmMaxPerRank = { + 0, // [0] Broadcast -- not used + 0, // [1] Reduce -- not used + 0, // [2] AllGather -- VMM off + 0, // [3] ReduceScatter -- ddaVmmMax total: Ring LL128 beats VMM 1.3-1.5x at + // 32M/64M on 16/32 ranks (DPX) + 0, // [4] AllReduce -- ddaVmmMax total (see ddaArVmmFirst*) + 0, // [5] SendRecv -- not used + 0, // [6] Send -- not used + 0, // [7] Recv -- not used + 0, // [8] AlltoAll -- VMM off + }, + .rankScaledMinRanks = 8, + .rankScaledMaxBytes = 256ULL*1024*1024, + // From 32 ranks, AllReduce VMM beats LL two-shot between 2M and nRanks x 256K; above that LL + // two-shot wins 1.2-1.3x (16M on 32 ranks) (DPX). + .ddaArVmmFirstMinRanks = 32, + .ddaArVmmFirstMinBytes = 2ULL*1024*1024, + .ddaArVmmFirstMaxPerRank = 256ULL*1024, + // VMM peaks at 128 / 128 / 64 / 32 blocks on 8 / 16 / 32 / 64 ranks. + .ddaVmmBlockBudget = 2048, + .ddaVmmMinBlocks = 32, + .ddaVmmMaxBlocks = 128, + // Ring / Tree crossovers, same sweep (AllReduce LL128: AICOMRCCL-2874/2875). + .arRingLL128MaxPerRank = 16ULL*1024*1024, // Ring Simple beats LL128 1.1-1.3x from 256M/512M/1G + // at 8/16/32 (DPX) + .arTreeLL128Unlimited = true, // Tree LL128 beats Tree Simple 1.3-2.9x at 32M-16G + .arTreeMaxPerRank = 1ULL*1024*1024, // Tree beats Ring 1.4-2.3x to 8M/16M/32M/64M + .arTreeLLMaxPerRank = 256ULL*1024, // Tree LL beats LL128 1.3-2.1x to 2M/4M/8M at 8/16/32 + // (fp8, DPX) + .agRingLL128MaxPerRank = 16ULL*1024*1024, // Ring Simple beats LL128 1.2-1.4x from 256M/1G/2G + // at 8/16/32 (DPX) + .rsRingLL128MinPerRank = 1ULL*1024*1024, // Ring LL beats LL128 1.1-1.2x to 1M per rank + // (fp8 at 16/32 ranks, DPX) + .rsRingLL128MaxPerRank = 8ULL*1024*1024, // 2/4-byte types: Simple wins past 8M per rank + .rsRingLL128MaxPerRank1B = 16ULL*1024*1024, // fp8: Ring Simple beats LL128 1.2-1.3x from 32M per + // rank (DPX) + .rsRingSimpleCappedMaxPerRank = 32ULL*1024*1024,// 128 channels beat 192 by 1.1-1.3x in (8M, 32M] + .rsRingSimpleCappedChannels = 128, + // CE-Scratch trails the DDA / Ring / Direct kernels it displaces between the DDA + // ceiling and the scratch size (nRanks x 64 MiB) (AICOMRCCL-2876). + .forceCeScratch = false, }; // gfx950: DDA-IPC cap is 128 MiB for AR/AG/RS and 4 MiB for AlltoAll. No fabric LL/LL128. @@ -205,6 +497,7 @@ static const rcclArchThresholds rcclArchThresholds_gfx950 = { kThreshUnlimited, // [8] AlltoAll -- not used }, .symMinR2 = {0, 0, 0, 0, 0, 0, 0, 0, 0}, + .forceCeScratch = true, }; // gfx942: DDA-IPC cap is 8 MiB for AR/AG/RS and 4 MiB for AlltoAll. No fabric LL/LL128. @@ -240,11 +533,14 @@ static const rcclArchThresholds rcclArchThresholds_gfx942 = { kThreshUnlimited, // [8] AlltoAll -- not used }, .symMinR2 = {0, 0, 0, 0, 0, 0, 0, 0, 0}, + .forceCeScratch = true, }; -const rcclArchThresholds* rcclGetArchThresholds(const char* gcn) { +const rcclArchThresholds* rcclGetArchThresholds(const char* gcn, rcclTuningPartition partition) { if (gcn == nullptr) return nullptr; - if (IsArchMatch(gcn, "gfx1250")) return &rcclArchThresholds_gfx1250; + if (IsArchMatch(gcn, "gfx1250")) { + return partition == RCCL_TUNING_PARTITION_DPX ? &rcclArchThresholds_gfx1250_dpx : &rcclArchThresholds_gfx1250; + } if (IsArchMatch(gcn, "gfx950")) return &rcclArchThresholds_gfx950; if (IsArchMatch(gcn, "gfx942")) return &rcclArchThresholds_gfx942; return nullptr; diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 82696a805fe3..aa556fd4c30b 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -153,14 +153,72 @@ int32_t rcclGetProtoForGfx110x(ncclFunc_t collectiveFunc, size_t sizePerRank) { return returnVal; } +namespace { +inline const rcclArchThresholds* extAlgoArchTable(const ncclComm* comm); +} + +// True when the user set `name` to a non-empty value. +static bool rcclEnvIsSet(const char* name) { + const char* v = ncclGetEnv(name); + return v != nullptr && v[0] != '\0'; +} + +// The user restricted the protocol set via NCCL_PROTO. Cached like rcclNcclAlgoEnvIsSet(). +static bool rcclNcclProtoEnvIsSet() { + static int cached = -1; + if (cached == -1) cached = getenv("NCCL_PROTO") != nullptr ? 1 : 0; + return cached == 1; +} + +// The arch table's Ring / Tree crossovers come from single-node (one clique) sweeps, +// so they apply only when nNodes == 1; the rank-scaled ones also need +// rankScaledMinRanks ranks. Null when they do not apply. +static const rcclArchThresholds* rcclRingTreeTable(const struct ncclComm* comm, bool rankScaled) { + if (comm->nNodes != 1) return nullptr; + const rcclArchThresholds* t = extAlgoArchTable(comm); + if (t == nullptr) return nullptr; + if (rankScaled && (t->rankScaledMinRanks <= 0 || comm->nRanks < t->rankScaledMinRanks)) return nullptr; + return t; +} + +// Arch-table replacements for the tuning model's LL / LL128 ranges (rank-scaled entries). +static void rcclArchProtoRanges(const struct ncclComm* comm, const struct ncclTaskColl* info, uint64_t& llMax, + uint64_t& ll128Min, uint64_t& ll128Max) { + const rcclArchThresholds* t = rcclRingTreeTable(comm, /*rankScaled=*/true); + if (t == nullptr) return; + const uint64_t nRanks = (uint64_t)comm->nRanks; + if (info->func == ncclFuncAllReduce && info->algorithm == NCCL_ALGO_TREE && t->arTreeLLMaxPerRank != 0) { + llMax = std::max(llMax, nRanks * t->arTreeLLMaxPerRank); + ll128Min = std::max(ll128Min, nRanks * t->arTreeLLMaxPerRank); + } else if (info->func == ncclFuncAllGather && t->agRingLL128MaxPerRank != 0) { + ll128Max = std::max(ll128Max, t->agRingLL128MaxPerRank); + } else if (info->func == ncclFuncReduceScatter && t->rsRingLL128MaxPerRank != 0) { + // The table's band replaces the model's outright (it is narrower at the bottom). + ll128Min = t->rsRingLL128MinPerRank; + ll128Max = (ncclTypeSize(info->datatype) == 1 && t->rsRingLL128MaxPerRank1B != 0) ? t->rsRingLL128MaxPerRank1B + : t->rsRingLL128MaxPerRank; + } +} + +void rcclUpdateCollectiveAlgorithm(struct ncclComm* comm, size_t nBytes, struct ncclTaskColl* info, + float table[][NCCL_NUM_PROTOCOLS]) { + // With NCCL_PROTO set, rcclUpdateCollectiveProtocol leaves the protocol alone, so a + // switch to Tree would keep the protocol picked for Ring. + if (rcclNcclAlgoEnvIsSet() || rcclNcclProtoEnvIsSet()) return; + if (info->func != ncclFuncAllReduce || info->algorithm != NCCL_ALGO_RING) return; + const rcclArchThresholds* t = rcclRingTreeTable(comm, /*rankScaled=*/true); + if (t == nullptr || t->arTreeMaxPerRank == 0 || nBytes > (uint64_t)comm->nRanks * t->arTreeMaxPerRank) return; + // rcclUpdateCollectiveProtocol then picks Tree LL or LL128, so both must be on offer. + if (table[NCCL_ALGO_TREE][NCCL_PROTO_LL] == NCCL_ALGO_PROTO_IGNORE || + table[NCCL_ALGO_TREE][NCCL_PROTO_LL128] == NCCL_ALGO_PROTO_IGNORE) + return; + info->algorithm = NCCL_ALGO_TREE; +} + void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, struct ncclTaskColl* info) { // Honor user input for protocol choice - static int userProtocolInput = -2; + const bool userProtocolInput = rcclNcclProtoEnvIsSet(); size_t sizePerRank = rcclGetSizePerRank(info->func, nBytes, comm->nRanks); - if (userProtocolInput == -2) { - const char* protoStr = getenv("NCCL_PROTO"); - userProtocolInput = !protoStr ? 0 : 1; - } if (!userProtocolInput && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950") && comm->nNodes == 1 && (info->func == ncclFuncAllGather) && sizePerRank <= 88448) { @@ -210,6 +268,7 @@ void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, s // Only override model choices if min/max cutoff points are set in the tuning models if ((ll128Max != RCCL_LL_LIMITS_UNDEFINED) || (llMax != RCCL_LL_LIMITS_UNDEFINED)) { + rcclArchProtoRanges(comm, info, llMax, ll128Min, ll128Max); // Keep it simple unless otherwise required info->protocol = NCCL_PROTO_SIMPLE; if (sizePerRank <= llMax && sizePerRank > llMin) { @@ -223,6 +282,14 @@ void rcclUpdateCollectiveProtocol(struct ncclComm* comm, size_t const& nBytes, s ll128Max *= 3.8; // Scale max message size for n > 2 since Tree has special behavior at 2 nodes } // ll128Max += (log2i(comm->nNodes) - 1) * comm->minMaxLLRange[tunableIndex][NCCL_PROTO_LL128][RCCL_PROTOCOL_FACTOR_IDX]; + // Arch-table AllReduce LL128 ceilings (gfx1250: AICOMRCCL-2874/2875). + if (const rcclArchThresholds* t = rcclRingTreeTable(comm, /*rankScaled=*/false)) { + if (info->algorithm == NCCL_ALGO_TREE && t->arTreeLL128Unlimited) { + ll128Max = UINT64_MAX; + } else if (info->algorithm != NCCL_ALGO_TREE && t->arRingLL128MaxPerRank != 0) { + ll128Max = std::max(ll128Max, (uint64_t)comm->nRanks * t->arRingLL128MaxPerRank); + } + } } if (sizePerRank <= ll128Max && sizePerRank > ll128Min) { info->protocol = NCCL_PROTO_LL128; @@ -266,9 +333,21 @@ extern int64_t ncclParamMaxNchannels(); extern int64_t rcclParamForceCe(); RCCL_PARAM(ChannelTuningEnable, "CHANNEL_TUNING_ENABLE", 1); +void rcclUpdateCollectiveChannels(struct ncclComm* comm, size_t nBytes, const struct ncclTaskColl* info, int& nc) { + if (info->func != ncclFuncReduceScatter || info->algorithm != NCCL_ALGO_RING || info->protocol != NCCL_PROTO_SIMPLE) + return; + // A user channel count (NCCL_MIN/MAX_NCHANNELS) wins. + if (ncclParamMaxNchannels() > 0 || ncclParamMinNchannels() > 0) return; + const rcclArchThresholds* t = rcclRingTreeTable(comm, /*rankScaled=*/true); + if (t == nullptr || t->rsRingSimpleCappedChannels <= 0) return; + const size_t perRank = nBytes / (size_t)comm->nRanks; + if (perRank > t->rsRingLL128MaxPerRank && perRank <= t->rsRingSimpleCappedMaxPerRank) + nc = std::min(nc, t->rsRingSimpleCappedChannels); +} + ncclResult_t rcclOverrideChannels(struct ncclComm* comm, ncclFunc_t coll, size_t nBytes, int& nc) { if (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx1250")) { - INFO(NCCL_TUNING, "RCCL Channel Tuning not applied for gfx1250 (pending sweep data, AICOMRCCL-1756)"); + INFO(NCCL_TUNING, "RCCL channel thresholds not applied for gfx1250 (see rcclUpdateCollectiveChannels)"); return ncclSuccess; } const bool isGfx_110x_120x = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx110") || @@ -769,7 +848,7 @@ ncclResult_t rcclGetProtocolName(int protocol, const char** protocolName) { namespace { // This comm's arch table, or NULL when the arch has no DDA tuning at all. // comm->archThresholds is populated at init; the lookup covers comms assembled -// by hand (unit tests) so they resolve exactly like production ones. +// by hand (unit tests) and resolves them like an SPX comm. inline const rcclArchThresholds* extAlgoArchTable(const ncclComm* comm) { if (comm == nullptr) return nullptr; if (rcclParamIgnoreArchTable() && !IsArchMatch(comm->archName, "gfx942") @@ -857,6 +936,20 @@ bool rcclForceCeAllReduceEnabled(const ncclComm* comm) { return rcclForceCeAllReduceEnabledDef(rcclCeAllReduceArchDefault(comm)); } +// Whether CE-Scratch is forced for unregistered AllGather / AlltoAll / Gather / +// Scatter. A user-set RCCL_FORCE_CE decides (0 = never, otherwise always). Left +// at its default, the comm-level NCCL_CTA_POLICY_ZERO forces it, and otherwise +// the arch table's forceCeScratch decides. +bool rcclForceCeScratchEnabled(const ncclComm* comm) { + static const bool envSet = rcclEnvIsSet("RCCL_FORCE_CE"); + const int64_t param = rcclParamForceCe(); + if (param == 0) return false; + if (envSet) return true; + if (comm->config.CTAPolicy & NCCL_CTA_POLICY_ZERO) return true; + const rcclArchThresholds* table = extAlgoArchTable(comm); + return table == nullptr || table->forceCeScratch; +} + inline size_t rcclCeNonRegMinTab(const rcclArchThresholds* table, ncclFunc_t func) { if (table == nullptr) return 0; return (size_t)func < RCCL_DDA_FUNC_COUNT ? table->ceNonRegMin[(size_t)func] : 0; @@ -881,14 +974,45 @@ bool rcclAllGatherCeRegisteredWindow(const ncclComm* comm, size_t totalBytes, } -inline size_t rcclDdaLLThresholdTab(const rcclArchThresholds* table, ncclFunc_t func) { +// True when this arch scales its DDA tier caps by rank count and the comm is big enough. +inline bool rcclDdaRankScaled(const rcclArchThresholds* table, int nRanks) { + return table != nullptr && table->rankScaledMinRanks > 0 && nRanks >= table->rankScaledMinRanks; +} + +// Table cap for one DDA tier: nRanks * the per-rank entry (clipped to +// rankScaledMaxBytes) on rank-scaled comms when that entry is set, else the total. +inline size_t ddaTierCapFromTable(const rcclArchThresholds* table, const size_t* totals, const size_t* perRank, + ncclFunc_t func, int nRanks) { + const size_t share = funcThresholdFromTable(perRank, func); + if (share != 0 && rcclDdaRankScaled(table, nRanks)) { + return std::min((size_t)nRanks * share, table->rankScaledMaxBytes); + } + return funcThresholdFromTable(totals, func); +} + +inline size_t rcclDdaLLThresholdTab(const rcclArchThresholds* table, ncclFunc_t func, int nRanks) { size_t threshold; if (ddaThresholdFromEnv(rcclParamDdaLLThreshold(), &threshold)) return threshold; if (table == nullptr) return kDdaLLBaseDefault; - return funcThresholdFromTable(table->ddaLLMax, func); + return ddaTierCapFromTable(table, table->ddaLLMax, table->ddaLLMaxPerRank, func, nRanks); } size_t rcclDdaLLThreshold(const ncclComm* comm, ncclFunc_t func) { - return rcclDdaLLThresholdTab(extAlgoArchTable(comm), func); + return rcclDdaLLThresholdTab(extAlgoArchTable(comm), func, comm != nullptr ? comm->nRanks : 0); +} + +// RCCL_DDA_LL_TWOSHOT_THRESHOLD wins when the user sets it. On rank-scaled comms +// whose table scales the AllReduce LL tier, two-shot runs to that tier's cap; +// otherwise the env var's default applies. +inline size_t rcclDdaLLTwoShotThresholdTab(const rcclArchThresholds* table, int nRanks) { + static const bool envSet = rcclEnvIsSet("RCCL_DDA_LL_TWOSHOT_THRESHOLD"); + if (!envSet && rcclDdaRankScaled(table, nRanks) && + funcThresholdFromTable(table->ddaLLMaxPerRank, ncclFuncAllReduce) != 0) { + return rcclDdaLLThresholdTab(table, ncclFuncAllReduce, nRanks); + } + return (size_t)rcclParamDdaLLTwoShotThreshold(); +} +size_t rcclDdaLLTwoShotThreshold(const ncclComm* comm) { + return rcclDdaLLTwoShotThresholdTab(extAlgoArchTable(comm), comm != nullptr ? comm->nRanks : 0); } // Returns the LL128 size cap for this collective, or 0 if the tier is disabled. @@ -898,53 +1022,90 @@ size_t rcclDdaLLThreshold(const ncclComm* comm, ncclFunc_t func) { // 0: explicitly disabled for all collectives. // >=1: forced on; uses table value when available, else kDdaLL128BaseDefault. // Callers check "> 0" to decide enablement, so no separate Active() helper needed. -inline size_t rcclDdaLL128ThresholdTab(const rcclArchThresholds* table, ncclFunc_t func) { +inline size_t rcclDdaLL128ThresholdTab(const rcclArchThresholds* table, ncclFunc_t func, int nRanks) { size_t threshold; if (ddaThresholdFromEnv(rcclParamDdaLL128Threshold(), &threshold)) return threshold; const int64_t param = rcclParamDdaLL128(); if (param == 0) return 0; if (param == kDdaThresholdUnset) { if (table == nullptr) return 0; - return funcThresholdFromTable(table->ddaLL128Max, func); + return ddaTierCapFromTable(table, table->ddaLL128Max, table->ddaLL128MaxPerRank, func, nRanks); } // param > 0: user forced on if (table == nullptr) return kDdaLL128BaseDefault; - return funcThresholdFromTable(table->ddaLL128Max, func); + return ddaTierCapFromTable(table, table->ddaLL128Max, table->ddaLL128MaxPerRank, func, nRanks); } size_t rcclDdaLL128Threshold(const ncclComm* comm, ncclFunc_t func) { - return rcclDdaLL128ThresholdTab(extAlgoArchTable(comm), func); + return rcclDdaLL128ThresholdTab(extAlgoArchTable(comm), func, comm != nullptr ? comm->nRanks : 0); } -inline size_t rcclDdaVmmThresholdTab(const rcclArchThresholds* table, ncclFunc_t func) { +// RCCL_DDA_LL128_TWOSHOT_THRESHOLD wins when the user sets it. When the table +// runs an AllReduce LL128 tier, two-shot runs to that tier's cap (as LL two-shot +// does to the LL tier's); otherwise the env var's default applies. +inline size_t rcclDdaLL128TwoShotThresholdTab(const rcclArchThresholds* table, int nRanks) { + static const bool envSet = rcclEnvIsSet("RCCL_DDA_LL128_TWOSHOT_THRESHOLD"); + if (!envSet && table != nullptr && funcThresholdFromTable(table->ddaLL128Max, ncclFuncAllReduce) != 0) { + return rcclDdaLL128ThresholdTab(table, ncclFuncAllReduce, nRanks); + } + return (size_t)rcclParamDdaLL128TwoShotThreshold(); +} +size_t rcclDdaLL128TwoShotThreshold(const ncclComm* comm) { + return rcclDdaLL128TwoShotThresholdTab(extAlgoArchTable(comm), comm != nullptr ? comm->nRanks : 0); +} + +inline size_t rcclDdaVmmThresholdTab(const rcclArchThresholds* table, ncclFunc_t func, int nRanks) { size_t threshold; if (ddaThresholdFromEnv(rcclParamDdaThreshold(), &threshold)) return threshold; if (table == nullptr) return kDdaVmmBaseDefault; - return funcThresholdFromTable(table->ddaVmmMax, func); + return ddaTierCapFromTable(table, table->ddaVmmMax, table->ddaVmmMaxPerRank, func, nRanks); } size_t rcclDdaVmmThreshold(const ncclComm* comm, ncclFunc_t func) { - return rcclDdaVmmThresholdTab(extAlgoArchTable(comm), func); + return rcclDdaVmmThresholdTab(extAlgoArchTable(comm), func, comm != nullptr ? comm->nRanks : 0); +} + +// AllReduce DDA VMM ahead of the LL tier (see rcclArchThresholds::ddaArVmmFirst*). +// Measured on unregistered buffers only; registered ones keep the tier order and +// the R2 cap. Env thresholds turn it off so RCCL_DDA_*_THRESHOLD keep their meaning. +inline bool rcclDdaArVmmFirstTab(const rcclArchThresholds* table, int nRanks, size_t msgBytes, + ncclSymRegType_t winRegType) { + if (table == nullptr || table->ddaArVmmFirstMinRanks <= 0 || nRanks < table->ddaArVmmFirstMinRanks) return false; + if (winRegType != ncclSymSendNonregRecvNonreg) return false; + size_t unused; + if (ddaThresholdFromEnv(rcclParamDdaThreshold(), &unused) || ddaThresholdFromEnv(rcclParamDdaLLThreshold(), &unused)) + return false; + return msgBytes > table->ddaArVmmFirstMinBytes && msgBytes <= (size_t)nRanks * table->ddaArVmmFirstMaxPerRank; +} + +int rcclDdaFabricVmmBlocks(const ncclComm* comm) { + if (comm == nullptr) return 0; + const int cap = comm->ddaFabricMaxBlocks; + const rcclArchThresholds* table = extAlgoArchTable(comm); + if (!rcclDdaRankScaled(table, comm->nRanks) || table->ddaVmmBlockBudget <= 0) return cap; + const int want = + std::min(table->ddaVmmMaxBlocks, std::max(table->ddaVmmMinBlocks, table->ddaVmmBlockBudget / comm->nRanks)); + return std::min(cap, want); } // Widest window any enabled DDA tier can serve for this collective. The entry // gate only decides whether the per-tier checks run, so a tier tuned off (0) // must not hide the tiers that are still on: gfx1250 ReduceScatter has // ddaVmmMax = 0 with non-zero LL/LL128 caps, and ddaLL128Max[AllReduce] -// (32 MiB) sits above ddaVmmMax[AllReduce] (16 MiB). RCCL_DDA_LL / +// (256 MiB) sits above ddaVmmMax[AllReduce] (16 MiB). RCCL_DDA_LL / // RCCL_DDA_LL128 = 0 drop their tier so disabling one cannot widen the gate. // Each tier still applies its own cap at the call site, including the VMM/IPC // branch, which checks rcclDdaVmmThreshold() there. // ddaVmmMaxGraph is also folded in: it is 0 for all collectives except AR // (256 MiB on gfx1250), so it only widens the gate for graph-mode AR calls // and is a no-op for all other collectives. -inline size_t rcclDdaEntryThresholdTab(const rcclArchThresholds* table, ncclFunc_t func) { - size_t cap = rcclDdaVmmThresholdTab(table, func); - if (rcclParamDdaLL()) cap = std::max(cap, rcclDdaLLThresholdTab(table, func)); - cap = std::max(cap, rcclDdaLL128ThresholdTab(table, func)); +inline size_t rcclDdaEntryThresholdTab(const rcclArchThresholds* table, ncclFunc_t func, int nRanks) { + size_t cap = rcclDdaVmmThresholdTab(table, func, nRanks); + if (rcclParamDdaLL()) cap = std::max(cap, rcclDdaLLThresholdTab(table, func, nRanks)); + cap = std::max(cap, rcclDdaLL128ThresholdTab(table, func, nRanks)); if (table != nullptr) cap = std::max(cap, funcThresholdFromTable(table->ddaVmmMaxGraph, func)); return cap; } size_t rcclDdaEntryThreshold(const ncclComm* comm, ncclFunc_t func) { - return rcclDdaEntryThresholdTab(extAlgoArchTable(comm), func); + return rcclDdaEntryThresholdTab(extAlgoArchTable(comm), func, comm != nullptr ? comm->nRanks : 0); } size_t rcclDdaScratchPayloadCap(const ncclComm* comm) { @@ -971,7 +1132,8 @@ size_t rcclDdaScratchPayloadCap(const ncclComm* comm) { } // All table entries are total message bytes; the selector compares totalBytes - // directly (rccl_wrap.cc rcclSelectReduceScatter), so no per-rank scaling here. + // directly (rccl_wrap.cc rcclSelectReduceScatter). Rank-scaled caps are clipped to +// rankScaledMaxBytes, which the table keeps within these totals. for (int i = 0; i < RCCL_DDA_FUNC_COUNT; ++i) { bump(table->ddaLLMax[i]); bump(table->ddaLL128Max[i]); @@ -998,7 +1160,7 @@ size_t rcclDdaScratchPayloadCap(const ncclComm* comm) { // graph-mode override, that cap wins. Graph-mode is checked first. // Env var (RCCL_DDA_THRESHOLD) always wins over all context variants. inline size_t rcclDdaVmmThresholdCtxTab(const rcclArchThresholds* table, ncclFunc_t func, - ncclSymRegType_t winRegType, bool graphMode) { + ncclSymRegType_t winRegType, bool graphMode, int nRanks) { size_t threshold; if (ddaThresholdFromEnv(rcclParamDdaThreshold(), &threshold)) return threshold; if (table == nullptr) return kDdaVmmBaseDefault; @@ -1012,7 +1174,7 @@ inline size_t rcclDdaVmmThresholdCtxTab(const rcclArchThresholds* table, ncclFun size_t r2Cap = funcThresholdFromTable(table->ddaVmmMaxR2, func); if (r2Cap != 0) return r2Cap; } - return funcThresholdFromTable(table->ddaVmmMax, func); + return ddaTierCapFromTable(table, table->ddaVmmMax, table->ddaVmmMaxPerRank, func, nRanks); } bool rcclDdaEnabled(const ncclComm* comm, size_t totalBytes, size_t threshold, @@ -1438,12 +1600,20 @@ ncclResult_t rcclSelectAllReduce(struct ncclComm* comm, const void* sendbuff, vo } #endif const bool ddaFabricArch1250 = IsArchMatch(comm->archName, "gfx1250"); - const size_t arDdaVmmMax = rcclDdaVmmThresholdCtxTab(archTable, ncclFuncAllReduce, winRegType, ceCapturing); + const size_t arDdaVmmMax = + rcclDdaVmmThresholdCtxTab(archTable, ncclFuncAllReduce, winRegType, ceCapturing, comm->nRanks); const bool arShouldTakeDda = rcclAllReduceShouldTakeDdaPath(comm, count, datatype, ddaSymEligible, ceAllReduceAllowed, query); if (arShouldTakeDda) { if (ddaFabricArch1250) { - const size_t arDdaLLMax = rcclDdaLLThresholdTab(archTable, ncclFuncAllReduce); - const size_t arDdaLL128Max = rcclDdaLL128ThresholdTab(archTable, ncclFuncAllReduce); + const size_t arDdaLLMax = rcclDdaLLThresholdTab(archTable, ncclFuncAllReduce, comm->nRanks); + const size_t arDdaLL128Max = rcclDdaLL128ThresholdTab(archTable, ncclFuncAllReduce, comm->nRanks); + // Large comms: VMM ahead of LL two-shot for mid sizes (ddaArVmmFirst*). + if (rcclDdaArVmmFirstTab(archTable, comm->nRanks, msgBytes, winRegType) && + ncclAllReduceDdaFabricEligible(comm, sendbuff, recvbuff, count, datatype, op)) { + decision->algo = RCCL_DDA_FABRIC_VMM; + decision->nMaxChannels = ncclAllReduceDdaFabricBlocks(comm, count, datatype); + return ncclSuccess; + } // Small-message fast lane: LL protocol (no GPU barrier). if (rcclParamDdaLL() && msgBytes <= arDdaLLMax && ncclAllReduceDdaFabricLLEligible(comm, sendbuff, recvbuff, count, datatype, op)) { @@ -1595,12 +1765,14 @@ ncclResult_t rcclSelectAllGather(struct ncclComm* comm, const void* sendbuff, vo // the CE-registered check so it loses to CE exactly as dispatch does // (taskAppend appends the CE task before ncclMakeSymmetricTaskList runs, so // symk never reclaims it), mirroring rcclSelectAllReduce. - const size_t agDdaVmmMax = rcclDdaVmmThresholdCtxTab(archTable, ncclFuncAllGather, winRegType, ceCapturing); - if (!symEligible && rcclDdaEnabled(comm, totalBytes, rcclDdaEntryThresholdTab(archTable, ncclFuncAllGather), query, "AG")) { + const size_t agDdaVmmMax = + rcclDdaVmmThresholdCtxTab(archTable, ncclFuncAllGather, winRegType, ceCapturing, comm->nRanks); + const size_t agDdaEntryMax = rcclDdaEntryThresholdTab(archTable, ncclFuncAllGather, comm->nRanks); + if (!symEligible && rcclDdaEnabled(comm, totalBytes, agDdaEntryMax, query, "AG")) { const bool agFabricArch = IsArchMatch(comm->archName, "gfx1250"); if (agFabricArch) { - const size_t agDdaLLMax = rcclDdaLLThresholdTab(archTable, ncclFuncAllGather); - const size_t agDdaLL128Max = rcclDdaLL128ThresholdTab(archTable, ncclFuncAllGather); + const size_t agDdaLLMax = rcclDdaLLThresholdTab(archTable, ncclFuncAllGather, comm->nRanks); + const size_t agDdaLL128Max = rcclDdaLL128ThresholdTab(archTable, ncclFuncAllGather, comm->nRanks); if (rcclParamDdaLL() && msgSize <= agDdaLLMax && ncclAllGatherDdaFabricLLEligible(comm, sendbuff, recvbuff, sendcount, datatype)) { decision->algo = RCCL_DDA_FABRIC_LL; @@ -1695,15 +1867,15 @@ ncclResult_t rcclSelectAllGather(struct ncclComm* comm, const void* sendbuff, vo const bool agCeNonRegWindow = agCeNonRegMax > 0 && totalBytes >= agCeNonRegMin && totalBytes <= agCeNonRegMax; - if ((rcclParamForceCe() || agCeNonRegWindow) && ceScratch && + if ((rcclForceCeScratchEnabled(comm) || agCeNonRegWindow) && ceScratch && winRegType == ncclSymSendNonregRecvNonreg && !hasSysmemSegment && comm->ddaScratch != nullptr && totalBytes <= (size_t)comm->ddaScratchBytes) { decision->algo = RCCL_CE_SCRATCH; return ncclSuccess; } - if (!query) INFO(NCCL_TUNING, "AG CE-scratch disqualified: ceScratch=%d hasSysmem=%d ddaScratch=%p totalBytes=%zu ddaScratchBytes=%zu forceCe=%d agCeNonRegWindow=%d", - (int)ceScratch, (int)hasSysmemSegment, comm->ddaScratch, totalBytes, (size_t)comm->ddaScratchBytes, (int)rcclParamForceCe(), (int)agCeNonRegWindow); + if (!query) INFO(NCCL_TUNING, "AG CE-scratch disqualified: ceScratch=%d hasSysmem=%d ddaScratch=%p totalBytes=%zu ddaScratchBytes=%zu forceCeScratch=%d agCeNonRegWindow=%d", + (int)ceScratch, (int)hasSysmemSegment, comm->ddaScratch, totalBytes, (size_t)comm->ddaScratchBytes, (int)rcclForceCeScratchEnabled(comm), (int)agCeNonRegWindow); // Branch #3: CE via registered symmetric windows. Taken when symk is not // eligible and the size is within ceRegMax, or when CTAPolicy=ZERO forces // CE for every size. @@ -1810,13 +1982,14 @@ ncclResult_t rcclSelectReduceScatter(struct ncclComm* comm, const void* sendbuff // (2) DDA fast paths. Symmetric wins when buffers are registered (-R 2); DDA // enters only when symk is unavailable. No Blocks helpers -> nMaxChannels 0. - const size_t rsDdaVmmMax = rcclDdaVmmThresholdCtxTab(archTable, ncclFuncReduceScatter, rsWinRegType, /*graphMode=*/false); - if (!symEligible && - rcclDdaEnabled(comm, totalBytes, rcclDdaEntryThresholdTab(archTable, ncclFuncReduceScatter), query, "RS")) { + const size_t rsDdaVmmMax = + rcclDdaVmmThresholdCtxTab(archTable, ncclFuncReduceScatter, rsWinRegType, /*graphMode=*/false, comm->nRanks); + const size_t rsDdaEntryMax = rcclDdaEntryThresholdTab(archTable, ncclFuncReduceScatter, comm->nRanks); + if (!symEligible && rcclDdaEnabled(comm, totalBytes, rsDdaEntryMax, query, "RS")) { const bool ddaFabricArch = IsArchMatch(comm->archName, "gfx1250"); if (ddaFabricArch) { - const size_t rsDdaLLMax = rcclDdaLLThresholdTab(archTable, ncclFuncReduceScatter); - const size_t rsDdaLL128Max = rcclDdaLL128ThresholdTab(archTable, ncclFuncReduceScatter); + const size_t rsDdaLLMax = rcclDdaLLThresholdTab(archTable, ncclFuncReduceScatter, comm->nRanks); + const size_t rsDdaLL128Max = rcclDdaLL128ThresholdTab(archTable, ncclFuncReduceScatter, comm->nRanks); if (rcclParamDdaLL() && totalBytes <= rsDdaLLMax && ncclReduceScatterDdaFabricLLEligible(comm, sendbuff, recvbuff, recvcount, datatype, op)) { decision->algo = RCCL_DDA_FABRIC_LL; @@ -1830,7 +2003,8 @@ ncclResult_t rcclSelectReduceScatter(struct ncclComm* comm, const void* sendbuff return ncclSuccess; } // The entry gate now admits the widest tier, so VMM re-checks its own cap. - // ddaVmmMax[RS] = 0 on gfx1250 keeps this tier off while LL/LL128 run. + // On gfx1250 SPX this tier runs only on rank-scaled comms (ddaVmmMaxPerRank[RS]); the + // DPX table's ddaVmmMax total sits under the LL tier except at 8 ranks. if (rsDdaVmmMax != 0 && totalBytes <= rsDdaVmmMax && ncclReduceScatterDdaFabricEligible(comm, sendbuff, recvbuff, recvcount, datatype, op)) { decision->algo = RCCL_DDA_FABRIC_VMM; @@ -1977,12 +2151,13 @@ ncclResult_t rcclSelectAlltoAll(struct ncclComm* comm, const void* sendbuff, voi // (3) DDA fast paths. gfx1250 uses fabric tiers; other archs use IPC. // Symmetric-registered buffers defer to the symmetric kernel; DDA gated on !a2aSymEligible. - const size_t a2aDdaMax = rcclDdaVmmThresholdTab(archTable, ncclFuncAlltoAll); - if (!a2aSymEligible && rcclDdaEnabled(comm, totalBytes, rcclDdaEntryThresholdTab(archTable, ncclFuncAlltoAll), query, "A2A")) { + const size_t a2aDdaMax = rcclDdaVmmThresholdTab(archTable, ncclFuncAlltoAll, comm->nRanks); + const size_t a2aDdaEntryMax = rcclDdaEntryThresholdTab(archTable, ncclFuncAlltoAll, comm->nRanks); + if (!a2aSymEligible && rcclDdaEnabled(comm, totalBytes, a2aDdaEntryMax, query, "A2A")) { const bool a2aFabricArch = IsArchMatch(comm->archName, "gfx1250"); if (a2aFabricArch) { - const size_t llThresh = rcclDdaLLThresholdTab(archTable, ncclFuncAlltoAll); - const size_t ll128Thresh = rcclDdaLL128ThresholdTab(archTable, ncclFuncAlltoAll); + const size_t llThresh = rcclDdaLLThresholdTab(archTable, ncclFuncAlltoAll, comm->nRanks); + const size_t ll128Thresh = rcclDdaLL128ThresholdTab(archTable, ncclFuncAlltoAll, comm->nRanks); if (rcclParamDdaLL() && llThresh > 0 && totalBytes <= llThresh && ncclAllToAllDdaFabricLLEligible(comm, sendbuff, recvbuff, count, datatype)) { decision->algo = RCCL_DDA_FABRIC_LL; @@ -2063,7 +2238,7 @@ ncclResult_t rcclSelectAlltoAll(struct ncclComm* comm, const void* sendbuff, voi // Guard a2aWinRegType to reject registered recv windows (they must not be routed through // ddaScratch; that would clobber the user recv buffer). Mirrors AllGather Branch #2 at // line ~1643 and the taskAppend funnel guard at enqueue.cc:3978. - if (rcclParamForceCe() && !a2aHasSysmem && + if (rcclForceCeScratchEnabled(comm) && !a2aHasSysmem && comm->ddaScratch != nullptr && totalBytes <= comm->ddaScratchBytes && ncclCeScratchAvailable(comm, ncclFuncAlltoAll, ncclDevSum, datatype, a2aWinRegType)) { decision->algo = RCCL_CE_SCRATCH; diff --git a/projects/rccl/test/DdaFabricEligibilityTests.cpp b/projects/rccl/test/DdaFabricEligibilityTests.cpp index 1dc48b878ffc..2a8215c65b2a 100644 --- a/projects/rccl/test/DdaFabricEligibilityTests.cpp +++ b/projects/rccl/test/DdaFabricEligibilityTests.cpp @@ -863,9 +863,9 @@ TEST_F(DdaFabricEligibilityTest, AllReduce_InvalidDatatypeDispatch) // size, with one-shot tested first. // // These run under the default parameter values -- DDA_LL=1, a 1 MiB -// DDA_LL_ONESHOT_THRESHOLD and a 16 MiB DDA_LL_TWOSHOT_THRESHOLD -- so sizes -// above 1 MiB and up to 16 MiB are claimed by the two-shot tier, which is what -// makes the tier split observable from here. They deliberately do not setenv: +// DDA_LL_ONESHOT_THRESHOLD and a 2 MiB DDA_LL_TWOSHOT_THRESHOLD, which gfx1250 +// raises to nRanks x 1 MiB from 8 ranks -- so sizes above 1 MiB are claimed by +// the two-shot tier, which is what makes the tier split observable from here. They deliberately do not setenv: // RCCL_PARAM caches on first read, and the NCCL_NO_CACHE opt-out is resolved once // per process, so a test that changed a threshold would depend on being the first // in the binary to touch it. Two guards stay out of reach as a result: @@ -1011,9 +1011,10 @@ TEST_F(DdaFabricEligibilityTest, AllReduceLL_TwoShotClaimsPastOneShotThreshold) } // 524288 floats is exactly the default 2 MiB DDA_LL_TWOSHOT_THRESHOLD, the top -// of the LL range, which the tier takes with <=. +// of the LL range on a comm below gfx1250's rank-scaled range, which the tier takes with <=. TEST_F(DdaFabricEligibilityTest, AllReduceLL_AtTwoShotThresholdEligible) { + mockComm_.comm.nRanks = 4; EXPECT_TRUE(ncclAllReduceDdaFabricLLEligible( mockComm_.get(), sendbuff_, recvbuff_, 524288, ncclFloat32, ncclSum)); } @@ -1021,14 +1022,31 @@ TEST_F(DdaFabricEligibilityTest, AllReduceLL_AtTwoShotThresholdEligible) // 128 bytes past the two-shot threshold: neither LL tier claims it and the message // falls through to LL128 / Simple. The count stays a multiple of 16 * nRanks so // the rejection is attributable to the threshold and not to the shard shape rules. +// Four ranks: below gfx1250's rank-scaled range, so the 2 MiB env default applies. TEST_F(DdaFabricEligibilityTest, AllReduceLL_PastBothThresholds) { + mockComm_.comm.nRanks = 4; // DDA_LL_TWOSHOT_THRESHOLD default 2 MiB; use 524320 float32 (~2.1 MiB) // so both one-shot (1 MiB) and two-shot (2 MiB) limits are exceeded. EXPECT_FALSE(ncclAllReduceDdaFabricLLEligible( mockComm_.get(), sendbuff_, recvbuff_, 524320, ncclFloat32, ncclSum)); } +// From 8 ranks, gfx1250 scales the AllReduce LL tier with the comm (ddaLLMaxPerRank, +// 1 MiB per rank), and two-shot runs to that cap: 8 MiB at 8 ranks, 16 MiB at 16. +TEST_F(DdaFabricEligibilityTest, AllReduceLL_TwoShotScalesWithRanks) +{ + for (int nRanks : {8, 16}) + { + mockComm_.comm.nRanks = nRanks; + const size_t capFloats = static_cast(nRanks) * 262144; // nRanks x 1 MiB + EXPECT_TRUE(ddaLLArTwoShotEligible( + mockComm_.get(), sendbuff_, recvbuff_, capFloats, ncclFloat32, ncclSum)) << nRanks; + EXPECT_FALSE(ddaLLArTwoShotEligible( + mockComm_.get(), sendbuff_, recvbuff_, capFloats + 4 * nRanks, ncclFloat32, ncclSum)) << nRanks; + } +} + // The gate above only reports the disjunction of the two tiers. These call each // per-variant predicate directly, so a tier can be pinned without the other // masking it, and assert the complementary tier to show which one owns the size. @@ -1222,12 +1240,13 @@ TEST_F(DdaFabricEligibilityTest, AllReduceLL128TwoShot_HalvedSlotStillReachesFur } // ncclAllReduceDdaFabricLL128Eligible threshold tests: at-cap and past-cap. -// These call the shape-eligibility function directly (which uses rcclParamDdaLL128OneShotThreshold / -// TwoShotThreshold, not the arch table), so they are independent of archName. -// The 32 MiB cap matches gfx1250 ddaLL128Max[AllReduce] from the arch table. +// These call the shape-eligibility function directly, which applies the one-shot and +// two-shot ceilings but not the tier's ddaLL128Max. 32 MiB was gfx1250's +// ddaLL128Max[AllReduce] when these were written; the probe size is kept as an +// in-range point. TEST_F(DdaFabricEligibilityTest, AllReduceLL128_AtCapEligible) { - // gfx1250 ddaLL128Max[AllReduce] = 32 MiB; 32 MiB / 4 bytes = 8388608 float32 + // 32 MiB / 4 bytes = 8388608 float32 EXPECT_TRUE(ncclAllReduceDdaFabricLL128Eligible( mockComm_.get(), sendbuff_, recvbuff_, 8388608, ncclFloat32, ncclSum)); } @@ -1239,6 +1258,37 @@ TEST_F(DdaFabricEligibilityTest, AllReduceLL128_PastCapRejected) mockComm_.get(), sendbuff_, recvbuff_, 8388610, ncclFloat32, ncclSum)); } +// gfx1250's two-shot runs to its AllReduce LL128 tier cap, 256 MiB; gfx950 has no +// LL128 tier and keeps RCCL_DDA_LL128_TWOSHOT_THRESHOLD's 64 MiB default. The fixture's +// scratch never limits the shard, so only the ceiling can reject. The cap + 128 bytes +// still leaves a 16-byte-aligned shard at 8 ranks, so the shard rule does not reject it. +TEST_F(DdaFabricEligibilityTest, AllReduceLL128TwoShot_CeilingFromArchTable) +{ + mockComm_.comm.nRanks = 8; + EXPECT_TRUE(ddaLL128ArTwoShotEligible( + mockComm_.get(), sendbuff_, recvbuff_, 67108864, ncclFloat32, ncclSum)); + EXPECT_FALSE(ddaLL128ArTwoShotEligible( + mockComm_.get(), sendbuff_, recvbuff_, 67108896, ncclFloat32, ncclSum)); + + mockComm_.comm.archName = const_cast("gfx950"); + EXPECT_TRUE(ddaLL128ArTwoShotEligible( + mockComm_.get(), sendbuff_, recvbuff_, 16777216, ncclFloat32, ncclSum)); + EXPECT_FALSE(ddaLL128ArTwoShotEligible( + mockComm_.get(), sendbuff_, recvbuff_, 16777248, ncclFloat32, ncclSum)); +} + +// The VMM AllReduce launch uses the rank-scaled block count: 128 blocks at 8 ranks and +// 64 at 32 for a message large enough to fill every block, from a 256-block init cap. +TEST_F(DdaFabricEligibilityTest, AllReduceVmm_BlocksFollowRankBudget) +{ + const size_t count = 64u * 1024u * 1024u; // 256 MiB of float32 + ASSERT_GE(mockComm_.comm.ddaFabricMaxBlocks, 128); + mockComm_.comm.nRanks = 8; + EXPECT_EQ(ncclAllReduceDdaFabricBlocks(mockComm_.get(), count, ncclFloat32), 128u); + mockComm_.comm.nRanks = 32; + EXPECT_EQ(ncclAllReduceDdaFabricBlocks(mockComm_.get(), count, ncclFloat32), 64u); +} + // --------------------------------------------------------------------------- // AllToAll // --------------------------------------------------------------------------- @@ -1490,16 +1540,21 @@ TEST_F(DdaFabricLL128EligibilityTest, AllGatherLL128_ScratchTooSmallForOneSlice) TEST_F(DdaFabricLL128EligibilityTest, AllGatherLL128_AtThresholdEligible) { - // fixture sets archName="gfx1250"; resolved threshold = ddaLL128Max[AG] = 64 MiB - constexpr size_t totalCap = 64u * 1024u * 1024u; - const size_t perRankBytes = totalCap / (size_t)mockComm_.comm.nRanks; - EXPECT_TRUE(ncclAllGatherDdaFabricLL128Eligible( - mockComm_.get(), sendbuff_, recvbuff_, perRankBytes / sizeof(float), ncclFloat32)); + // fixture sets archName="gfx1250": 128 MiB at 8 ranks, 256 MiB (clipped) at 16 + for (int nRanks : {8, 16}) + { + mockComm_.comm.nRanks = nRanks; + const size_t totalCap = rcclDdaLL128Threshold(mockComm_.get(), ncclFuncAllGather); + EXPECT_EQ(totalCap, (nRanks == 8 ? 128u : 256u) * 1024u * 1024u) << nRanks; + const size_t perRankBytes = totalCap / (size_t)nRanks; + EXPECT_TRUE(ncclAllGatherDdaFabricLL128Eligible( + mockComm_.get(), sendbuff_, recvbuff_, perRankBytes / sizeof(float), ncclFloat32)) << nRanks; + } } TEST_F(DdaFabricLL128EligibilityTest, AllGatherLL128_PastThresholdRejected) { - constexpr size_t totalCap = 64u * 1024u * 1024u; + const size_t totalCap = rcclDdaLL128Threshold(mockComm_.get(), ncclFuncAllGather); const size_t perRankBytes = totalCap / (size_t)mockComm_.comm.nRanks + 16; EXPECT_FALSE(ncclAllGatherDdaFabricLL128Eligible( mockComm_.get(), sendbuff_, recvbuff_, perRankBytes / sizeof(float), ncclFloat32)); diff --git a/projects/rccl/test/RcclWrapTests.cpp b/projects/rccl/test/RcclWrapTests.cpp index 36f0956cf9d4..2958f4852512 100644 --- a/projects/rccl/test/RcclWrapTests.cpp +++ b/projects/rccl/test/RcclWrapTests.cpp @@ -2641,7 +2641,7 @@ TEST(Rcclwrap, Gfx1250_CeEligible_StillTakesDda) { auto comm = std::make_unique(); InitDdaDecisionComm(*comm, "gfx1250", 8, 1, /*symmetricSupport=*/true); - size_t count = CountForBytes(16ull * 1024 * 1024, ncclFloat32); // CE-eligible (4-256 MiB window) and within DDA entry cap (32 MiB), divisible by 8 + size_t count = CountForBytes(16ull * 1024 * 1024, ncclFloat32); // CE-eligible (4-256 MiB window) and within the DDA entry cap, divisible by 8 EXPECT_TRUE(rcclAllReduceShouldTakeDdaPath(comm.get(), count, ncclFloat32, /*symEligible=*/false, /*ceAllReduceAllowed=*/true, /*query=*/false)); } @@ -3008,18 +3008,122 @@ TEST(RcclDdaTierThresholds, Gfx1250_NoEnv_UsesArchTable) InitDdaDecisionComm(*comm, "gfx1250", 8, 1, /*symmetricSupport=*/false); comm->archThresholds = rcclGetArchThresholds("gfx1250"); ASSERT_NE(comm->archThresholds, nullptr); - // LL and LL128 ceilings for AR are both 32 MiB on gfx1250; - // VMM ceiling is 16 MiB (above this, Ring/CE takes over). + // On an 8-rank gfx1250 comm the AR LL ceiling is rank-scaled to 8 MiB + // (1 MiB per rank) and the LL128 ceiling is 256 MiB; VMM ceiling is + // 16 MiB (above this, Ring/CE takes over). EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAllReduce), - 32ULL * 1024 * 1024); + 8ULL * 1024 * 1024); EXPECT_EQ(rcclDdaLL128Threshold(comm.get(), ncclFuncAllReduce), - 32ULL * 1024 * 1024); + 256ULL * 1024 * 1024); EXPECT_EQ(rcclDdaVmmThreshold(comm.get(), ncclFuncAllReduce), 16ULL * 1024 * 1024); }, {}); } +// The rank-scaled caps from the 8-64 rank H13 sweep (AICOMRCCL-1756): below 8 ranks +// the totals stay in force, from 8 ranks each tier grows with the comm and is +// clipped at 256 MiB. +TEST(RcclDdaTierThresholds, Gfx1250_RankScaledCaps) +{ + RUN_ISOLATED_TEST_WITH_ENV( + "Gfx1250_RankScaledCaps", + []() + { + constexpr size_t KiB = 1024, MiB = 1024 * KiB; + auto comm = std::make_unique(); + InitDdaDecisionComm(*comm, "gfx1250", 4, 1, /*symmetricSupport=*/false); + comm->archThresholds = rcclGetArchThresholds("gfx1250"); + ASSERT_NE(comm->archThresholds, nullptr); + // 4 ranks: totals. + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAllReduce), 32 * MiB); + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAllGather), 128 * KiB); + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncReduceScatter), 16 * MiB); + EXPECT_EQ(rcclDdaLL128Threshold(comm.get(), ncclFuncAllGather), 128 * MiB); + EXPECT_EQ(rcclDdaLL128Threshold(comm.get(), ncclFuncReduceScatter), 2 * MiB); + EXPECT_EQ(rcclDdaVmmThreshold(comm.get(), ncclFuncReduceScatter), 0ul); + for (int nRanks : {8, 16, 32, 64}) { + comm->nRanks = nRanks; + const size_t n = static_cast(nRanks); + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAllReduce), n * MiB) << nRanks; + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAllGather), n * 16 * KiB) << nRanks; + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncReduceScatter), n * 256 * KiB) << nRanks; + EXPECT_EQ(rcclDdaLL128Threshold(comm.get(), ncclFuncAllGather), std::min(n * 16 * MiB, 256 * MiB)) + << nRanks; + EXPECT_EQ(rcclDdaLL128Threshold(comm.get(), ncclFuncReduceScatter), n * 512 * KiB) << nRanks; + EXPECT_EQ(rcclDdaVmmThreshold(comm.get(), ncclFuncReduceScatter), n * 2 * MiB) << nRanks; + // Unscaled tiers keep their totals. + EXPECT_EQ(rcclDdaLL128Threshold(comm.get(), ncclFuncAllReduce), 256 * MiB) << nRanks; + EXPECT_EQ(rcclDdaVmmThreshold(comm.get(), ncclFuncAllReduce), 16 * MiB) << nRanks; + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAlltoAll), 64 * KiB) << nRanks; + } + }, + {}); +} + +// Only gfx1250 scales its caps with the comm: on gfx942/gfx950 a 64-rank comm resolves +// the same caps as an 8-rank one, and the two-shot ceilings keep the env defaults. +// On gfx1250 both two-shot ceilings follow their tier caps. +TEST(RcclDdaTierThresholds, RankScalingOnlyOnGfx1250) +{ + RUN_ISOLATED_TEST_WITH_ENV( + "RankScalingOnlyOnGfx1250", + []() + { + constexpr size_t MiB = 1024 * 1024; + auto comm = std::make_unique(); + for (const char* arch : {"gfx942", "gfx950"}) + { + InitDdaDecisionComm(*comm, arch, 8, 1, /*symmetricSupport=*/false); + comm->archThresholds = rcclGetArchThresholds(arch); + const size_t vmm8 = rcclDdaVmmThreshold(comm.get(), ncclFuncReduceScatter); + const size_t ll8 = rcclDdaLLThreshold(comm.get(), ncclFuncAllGather); + comm->nRanks = 64; + EXPECT_EQ(rcclDdaVmmThreshold(comm.get(), ncclFuncReduceScatter), vmm8) << arch; + EXPECT_EQ(rcclDdaLLThreshold(comm.get(), ncclFuncAllGather), ll8) << arch; + EXPECT_EQ(rcclDdaLLTwoShotThreshold(comm.get()), 2 * MiB) << arch; + EXPECT_EQ(rcclDdaLL128TwoShotThreshold(comm.get()), 64 * MiB) << arch; + } + InitDdaDecisionComm(*comm, "gfx1250", 16, 1, /*symmetricSupport=*/false); + comm->archThresholds = rcclGetArchThresholds("gfx1250"); + EXPECT_EQ(rcclDdaLLTwoShotThreshold(comm.get()), 16 * MiB); + EXPECT_EQ(rcclDdaLL128TwoShotThreshold(comm.get()), 256 * MiB); + }, + {}); +} + +// CE-Scratch is forced on gfx942/gfx950 as before and left to the kernel paths on +// gfx1250 (AICOMRCCL-2876). +TEST(RcclDdaTierThresholds, ForceCeScratchPerArch) +{ + EXPECT_FALSE(rcclGetArchThresholds("gfx1250")->forceCeScratch); + EXPECT_TRUE(rcclGetArchThresholds("gfx950")->forceCeScratch); + EXPECT_TRUE(rcclGetArchThresholds("gfx942")->forceCeScratch); +} + +// RCCL_TUNING_PARTITION_MODE: gfx1250 has its own DPX table, tuned on DPX nodes apart from +// the SPX one; the other archs return one table for every mode. +TEST(RcclDdaTierThresholds, TuningPartitionTablePerArch) +{ + const rcclArchThresholds* spx = rcclGetArchThresholds("gfx1250"); + const rcclArchThresholds* dpx = rcclGetArchThresholds("gfx1250", RCCL_TUNING_PARTITION_DPX); + ASSERT_NE(spx, nullptr); + ASSERT_NE(dpx, nullptr); + EXPECT_EQ(spx, rcclGetArchThresholds("gfx1250", RCCL_TUNING_PARTITION_SPX)); + EXPECT_NE(spx, dpx); + EXPECT_NE(0, std::memcmp(spx, dpx, sizeof(*spx))) << "the DPX table carries its own tuning"; + // A few of the DPX-only entries, so a DPX value copied back over SPX shows up here. + EXPECT_EQ(dpx->agRingLL128MaxPerRank, (size_t)16 * 1024 * 1024); + EXPECT_EQ(dpx->arTreeLLMaxPerRank, (size_t)256 * 1024); + EXPECT_EQ(dpx->ddaLLMax[ncclFuncAlltoAll], (size_t)2 * 1024 * 1024); + EXPECT_NE(spx->agRingLL128MaxPerRank, dpx->agRingLL128MaxPerRank); + for (const char* arch : {"gfx950", "gfx942"}) + { + EXPECT_EQ(rcclGetArchThresholds(arch), rcclGetArchThresholds(arch, RCCL_TUNING_PARTITION_DPX)) << arch; + } +} + + TEST(RcclDdaTierThresholds, Gfx1250_EnvOverridesArchTable) { RUN_ISOLATED_TEST_WITH_ENV( diff --git a/projects/rccl/test/host/CMakeLists.txt b/projects/rccl/test/host/CMakeLists.txt index 50e6482f437d..8006d5843755 100644 --- a/projects/rccl/test/host/CMakeLists.txt +++ b/projects/rccl/test/host/CMakeLists.txt @@ -936,11 +936,25 @@ endif() # --------------------------------------------------------------------------- # rccl-UnitTestsMicroCe # --------------------------------------------------------------------------- +# In-build, CE_BATCH_ASYNC_SUPPORTED arrives with rccl's own compile definitions when CE_ENABLED. +set(_MICRO_CE_BATCH_DEF "") +if(NOT BUILD_TESTS) + include(CheckCXXSymbolExists) + set(CMAKE_TRY_COMPILE_TARGET_TYPE STATIC_LIBRARY) + set(CMAKE_REQUIRED_DEFINITIONS -D__HIP_PLATFORM_AMD__) + check_cxx_symbol_exists(hipMemcpyBatchAsync "hip/hip_runtime_api.h" _MICRO_HIP_MEMCPY_BATCH_ASYNC) + unset(CMAKE_REQUIRED_DEFINITIONS) + set(CMAKE_TRY_COMPILE_TARGET_TYPE EXECUTABLE) + if(_MICRO_HIP_MEMCPY_BATCH_ASYNC) + set(_MICRO_CE_BATCH_DEF CE_BATCH_ASYNC_SUPPORTED) + endif() +endif() rccl_add_micro_binary(rccl-UnitTestsMicroCe SOURCES ${TEST_MICRO_CE_SOURCE_FILES} EXTRA_DEFS ${_MICRO_MODE_DEFS} ${_MICRO_AMDSMI_DEF} + ${_MICRO_CE_BATCH_DEF} RCCL_STUBS_OMIT_ncclCeFinalize # ce_coll.cc defines it CE_COLL_CC_PATH="${RCCL_HOST_HIPIFY_DIR}/src/ce_coll.cc" CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/../test_categories_micro.yaml") diff --git a/projects/rccl/test/host/MICROTEST_README.md b/projects/rccl/test/host/MICROTEST_README.md index ea7731b3d661..6d1c7e09ba75 100644 --- a/projects/rccl/test/host/MICROTEST_README.md +++ b/projects/rccl/test/host/MICROTEST_README.md @@ -193,8 +193,12 @@ export colliding non-`static` symbols; otherwise a unit needs its own binary: because `enqueue.cc` defines them itself. See `test_categories_micro_enqueue.yaml`. - **`rccl-UnitTestsMicroCe`**: `ce_coll.cc` (via `CE_COLL_CC_PATH`); suite - `CeCollMicrotest.*`. Its own binary because every other target links - `fakes/ce_fakes.cc`, which fakes the symbols this file defines. + `CeCollMicrotest.*` and `LaunchPaths/CeAlltoAllvSyncMicrotest.*`, both in + `ce-coll-test.cc`. Its own binary because every other target links + `fakes/ce_fakes.cc`, which fakes the symbols this file defines, and the + `rccl-UnitTestsMicroInit` cleanup-order oracle needs the fake `ncclCeFinalize` + in `fakes/nccl_stubs.cc`, which this target drops via + `RCCL_STUBS_OMIT_ncclCeFinalize`. - **`rccl-UnitTestsMicroSymKernels`** — the REAL `src/sym_kernels.cc` (via `SYM_KERNELS_CC_PATH`, from `sym-kernels-test.cc`), compiled together with the GENERATED `sym_kernels_host.cc` it calls into; suites `SymKernelMicrotest.*`, @@ -386,7 +390,7 @@ symbol. |---|---| | `src/algorithms/dda/*.cc` | `fakes/dda_fakes.cc` | | `src/bootstrap.cc` | `fakes/bootstrap_stubs.cc` | -| `src/ce_coll.cc` | `fakes/ce_fakes.cc` | +| `src/ce_coll.cc` (targets that do not compile the real file) | `fakes/ce_fakes.cc` | | `src/collectives.cc` | `fakes/collectives_fakes.cc` | | `src/dev_runtime.cc` (targets that do not compile the real file) | `fakes/dev_runtime_fakes.cc` | | `src/diagnostics.cc` (`ncclDiagChildRun*` external-tool runners) | `fakes/diagnostics_fakes.cc` | @@ -413,6 +417,7 @@ symbol. | `src/plugin/gin.cc`, `src/gin/gin_host.cc` (targets that do not compile the real file) | `fakes/gin_fakes.cc` | | `src/proxy.cc` | `fakes/proxy_fakes.cc` | | `src/ras/ras_param.cc` | `fakes/ras_param_fakes.cc` (the direct `ras-param-test.cc` inclusion uses renamed symbols) | +| `src/rccl_arch_thresholds.cc` | `fakes/rccl_arch_thresholds_fakes.cc` | | `src/rccl_wrap.cc`'s own public entry points (targets that don't compile the real file, e.g. `rccl-UnitTestsMicroEnqueue`) | `fakes/rccl_wrap_fakes.cc` | | `src/rccl_wrap.cc`'s dependencies (`rccl-UnitTestsMicro`, which compiles the real file and tests it directly) | `fakes/wrap_fakes.cc` | | `src/recorder.cc` | `fakes/recorder_fakes.cc` | @@ -773,7 +778,7 @@ above (`./install.sh -t`, wired via `add_subdirectory(host)`), the same file can be configured **directly** to build every host binary — `rccl-HostUnitTests`, `rccl-UnitTestsMicro`, `rccl-UnitTestsMicroWarpSpeed`, `rccl-UnitTestsMicroInit[-uncached|-faultinj]`, `rccl-UnitTestsMicroEnqueue[-devlinker]`, -`rccl-UnitTestsMicroGinHost, `rccl-UnitTestsMicroDiagnostics`, +`rccl-UnitTestsMicroGinHost`, `rccl-UnitTestsMicroCe`, `rccl-UnitTestsMicroDiagnostics`, `rccl-UnitTestsMicroSymKernels` and `rccl-UnitTestsMicroTaskPrep` — **without configuring/building all of librccl**. It compiles just the tests + fakes + the hipified unit-under-test sources. @@ -815,6 +820,7 @@ cmake --build build -j"$(nproc)" ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/{p2p,ib_write_bw}.cc tests ./build/rccl-UnitTestsMicroGinHost # src/gin/gin_host.cc GIN_PROXY_NTHREADS ./build/rccl-UnitTestsMicroDiagnostics # src/diagnostics/p2p.cc tests +./build/rccl-UnitTestsMicroCe # src/ce_coll.cc tests ./build/rccl-HostUnitTests ``` diff --git a/projects/rccl/test/host/ce-coll-test.cc b/projects/rccl/test/host/ce-coll-test.cc index 49f7457a3cf9..3b87dd8b3dfc 100644 --- a/projects/rccl/test/host/ce-coll-test.cc +++ b/projects/rccl/test/host/ce-coll-test.cc @@ -9,15 +9,21 @@ #include +#include +#include #include +#include #include +#include #include #include +#include #include "../common/LogCapture.hpp" #include "ScopedHook.h" #include "fakes/allocator_fakes.h" #include "fakes/dev_runtime_fakes.h" +#include "fakes/env_fakes.h" #include "fakes/hip_fakes.h" #include "fakes/nccl_fakes.h" // g_loadParam, used by param_redirect.h #include "fakes/nccl_stubs.h" @@ -214,4 +220,454 @@ TEST_F(CeCollMicrotest, EnsureStaging_AlreadySetUp_ReturnsWithoutRegistering) { EXPECT_EQ(0, alloc_.calls); } +constexpr size_t kMiB = 1024 * 1024; +constexpr uintptr_t kSendBase = 0x100000000000; +constexpr uintptr_t kRecvBase = 0x200000000000; +constexpr uintptr_t kPeerWindowBase = 0x300000000000; +constexpr uintptr_t kPeerWindowStride = 0x10000000000; +constexpr unsigned long long kCaptureGraphId = 7; +constexpr int kBatchAsyncDriverVersion = ROCM_VER_7_12_0; + +enum class LaunchPath { kSingleStream, kMultiStream, kBatchAsync, kLegacyNullStream, kGraphCapture }; + +std::string LaunchPathName(const ::testing::TestParamInfo& info) { + switch (info.param) { + case LaunchPath::kSingleStream: + return "SingleStream"; + case LaunchPath::kMultiStream: + return "MultiStream"; + case LaunchPath::kBatchAsync: + return "BatchAsync"; + case LaunchPath::kLegacyNullStream: + return "LegacyNullStream"; + case LaunchPath::kGraphCapture: + return "GraphCapture"; + } + return "Unknown"; +} + +using SizeMatrix = std::vector>; + +// Row r sends totalBytes over peerCounts[r] consecutive destinations from itself; the first takes any remainder. +SizeMatrix EqualTotalsMatrix(const std::vector& peerCounts, size_t totalBytes) { + const int nRanks = static_cast(peerCounts.size()); + SizeMatrix send(nRanks, std::vector(nRanks, 0)); + for (int src = 0; src < nRanks; ++src) { + for (int j = 0; j < peerCounts[src]; ++j) { + send[src][(src + j) % nRanks] = totalBytes / peerCounts[src] + (j == 0 ? totalBytes % peerCounts[src] : 0); + } + } + return send; +} + +std::vector CyclingPeerCounts(int nRanks, int minPeers, int span) { + std::vector counts(nRanks); + for (int r = 0; r < nRanks; ++r) { + counts[r] = minPeers + r % span; + } + return counts; +} + +std::vector GatherSizes(const SizeMatrix& send) { + const int nRanks = static_cast(send.size()); + std::vector gathered(4 * static_cast(nRanks) * nRanks); + for (int rank = 0; rank < nRanks; ++rank) { + std::vector sendDispls(nRanks); + std::vector recvSizes(nRanks); + std::vector recvDispls(nRanks); + size_t sendOff = 0; + size_t recvOff = 0; + for (int peer = 0; peer < nRanks; ++peer) { + sendDispls[peer] = sendOff; + sendOff += send[rank][peer]; + recvSizes[peer] = send[peer][rank]; + recvDispls[peer] = recvOff; + recvOff += recvSizes[peer]; + } + ncclAlltoAllvPackLocalSizes(ncclAlltoAllvRankMetaBlock(gathered.data(), rank, nRanks), nRanks, + send[rank].data(), sendDispls.data(), recvSizes.data(), recvDispls.data()); + } + return gathered; +} + +struct RankRun { + uint32_t barriers = 0; + size_t copies = 0; + size_t bytes = 0; + size_t batchCopies = 0; + size_t copyStreamCopies = 0; + size_t copyCalls = 0; + int syncCalls = 0; + // Barriers already issued when the copy to each destination slot was enqueued; 0 for a slot with nothing to send. + std::vector slotEpochs; +}; + +constexpr uintptr_t kCopyStreamBase = 0x6000; + +class CeAlltoAllvSyncMicrotest : public ::testing::TestWithParam { + protected: + void SetUp() override { + ResetAllFakes(); +#ifndef CE_BATCH_ASYNC_SUPPORTED + if (GetParam() == LaunchPath::kBatchAsync) { + GTEST_SKIP() << "hip_runtime_api.h has no hipMemcpyBatchAsync, so ce_coll.cc builds without the batch path"; + } +#endif + g_hipStreamBatchMemOp = [this](hipStream_t, unsigned int, hipStreamBatchMemOpParams*, unsigned int) { + return ++current_.syncCalls == failSyncCall_ ? hipErrorInvalidValue : hipSuccess; + }; + g_hipMemcpyAsync = [this](void* dst, const void*, size_t bytes, hipMemcpyKind, hipStream_t stream) { + if (++current_.copyCalls == failCopyCall_) { + return hipErrorInvalidValue; + } + RecordCopy(dst, bytes); + const uintptr_t s = reinterpret_cast(stream); + if (s >= kCopyStreamBase && s < kCopyStreamBase + RCCL_CE_NUM_COPY_STREAMS) { + ++current_.copyStreamCopies; + } + return hipSuccess; + }; +#ifdef CE_BATCH_ASYNC_SUPPORTED + g_hipMemcpyBatchAsync = [this](void** dsts, void**, size_t* sizes, size_t count, hipMemcpyAttributes*, + size_t*, size_t, size_t*, hipStream_t) { + if (++current_.copyCalls == failCopyCall_) { + return hipErrorInvalidValue; + } + for (size_t i = 0; i < count; ++i) { + RecordCopy(dsts[i], sizes[i]); + } + current_.batchCopies += count; + return hipSuccess; + }; +#endif + g_hipEventRecord = [](hipEvent_t, hipStream_t) { return hipSuccess; }; + g_hipStreamWaitEvent = [](hipStream_t, hipEvent_t, unsigned int) { return hipSuccess; }; + g_devrGetLsaRankPtr = [](ncclComm* comm, ncclDevrWindow*, size_t offset, int lsaRank, void** outPtr) { + if (lsaRank < 0 || lsaRank >= comm->devrState.lsaSize) { + return ncclInvalidArgument; + } + *outPtr = reinterpret_cast(kPeerWindowBase + lsaRank * kPeerWindowStride + offset); + return ncclSuccess; + }; + const bool batchAsync = GetParam() == LaunchPath::kBatchAsync; + g_loadParam = [batchAsync](const char* env, int64_t deftVal) -> int64_t { + if (std::strcmp(env, "RCCL_CE_BATCH_ASYNC_ENABLE") == 0) { + return batchAsync ? 1 : 0; + } + return deftVal; + }; + if (batchAsync) { + ncclCudaDriverVersionCache = kBatchAsyncDriverVersion; + } + } + + void TearDown() override { ResetAllFakes(); } + + static void ResetAllFakes() { + ResetHipFakes(); + ResetDevRuntimeFakes(); + ResetNcclFakes(); + ResetNcclStubs(); + ResetEnvFakes(); + } + + void RecordCopy(void* dst, size_t bytes) { + ++current_.copies; + current_.bytes += bytes; + const uintptr_t addr = reinterpret_cast(dst); + const int nRanks = activeComm_->nRanks; + const int dstRank = + addr >= kPeerWindowBase ? static_cast((addr - kPeerWindowBase) / kPeerWindowStride) : activeComm_->rank; + current_.slotEpochs[(dstRank - activeComm_->rank + nRanks) % nRanks] = activeComm_->ceColl.ceSeqNum; + } + + // Runs ncclCeAlltoAllv once per simulated rank; every rank sees the same gathered size matrix, as in production. + std::vector Run(const SizeMatrix& send, uint32_t freq, ncclResult_t expected = ncclSuccess) { + const int nRanks = static_cast(send.size()); + std::vector gathered = GatherSizes(send); + std::vector lsaRankList(nRanks); + for (int r = 0; r < nRanks; ++r) { + lsaRankList[r] = r; + } + std::vector syncFlags(2 * nRanks, 0); + ncclDevrWindow syncWin = {}; + syncWin.userPtr = syncFlags.data(); + ncclDevrWindow recvWin = {}; + recvWin.userPtr = reinterpret_cast(kRecvBase); + + const LaunchPath path = GetParam(); + hipStream_t stream = path == LaunchPath::kLegacyNullStream ? nullptr : reinterpret_cast(0x5717); + std::vector runs; + for (int rank = 0; rank < nRanks; ++rank) { + auto comm = std::make_unique(); + comm->rank = rank; + comm->nRanks = nRanks; + comm->devrState.lsaSize = nRanks; + comm->devrState.lsaSelf = rank; + comm->devrState.lsaRankList = lsaRankList.data(); + comm->symkState.hasLsaMultimem = false; + comm->planner.capturingGraph.graphId = path == LaunchPath::kGraphCapture ? kCaptureGraphId : ULLONG_MAX; + comm->ceColl.intraBatchSyncFreq = freq; + comm->ceColl.intraBatchSyncMsgThreshold = CE_COLL_INTRA_BATCH_SYNC_MSG_THRESHOLD; + comm->ceColl.baseUCSymReadyPtr = reinterpret_cast(syncFlags.data()); + comm->ceColl.baseUCSymComplPtr = reinterpret_cast(syncFlags.data() + nRanks); + comm->ceColl.ceSyncWin = &syncWin; + if (path == LaunchPath::kMultiStream) { + comm->ceColl.nCopyStreams = RCCL_CE_NUM_COPY_STREAMS; + for (int s = 0; s < RCCL_CE_NUM_COPY_STREAMS; ++s) { + comm->ceColl.copyStreams[s] = reinterpret_cast(kCopyStreamBase + s); + comm->ceColl.copyEvents[s] = reinterpret_cast(0x7000 + s); + } + } + + ncclCeCollArgs args = {}; + args.func = ncclFuncAlltoAllv; + args.sendBuff = reinterpret_cast(kSendBase); + args.recvBuff = reinterpret_cast(kRecvBase); + args.recvWin = &recvWin; + args.sizes = gathered.data(); + + current_ = RankRun{}; + current_.slotEpochs.assign(nRanks, 0); + activeComm_ = comm.get(); + EXPECT_EQ(ncclCeAlltoAllv(comm.get(), &args, stream), expected) << "rank " << rank; + activeComm_ = nullptr; + current_.barriers = comm->ceColl.ceSeqNum; + runs.push_back(current_); + } + return runs; + } + + static std::vector Barriers(const std::vector& runs) { + std::vector barriers; + for (const RankRun& run : runs) { + barriers.push_back(run.barriers); + } + return barriers; + } + + // Each rank copies exactly its nonzero row, slot s after the ready barrier and s / roundSlots round barriers. + static void ExpectRowsCopied(const SizeMatrix& send, const std::vector& runs, uint32_t roundSlots) { + const int nRanks = static_cast(send.size()); + for (int rank = 0; rank < nRanks; ++rank) { + size_t copies = 0; + size_t bytes = 0; + std::vector epochs(nRanks, 0); + for (int slot = 0; slot < nRanks; ++slot) { + const size_t chunk = send[rank][(rank + slot) % nRanks]; + if (chunk != 0) { + ++copies; + bytes += chunk; + epochs[slot] = 1 + slot / roundSlots; + } + } + EXPECT_EQ(runs[rank].copies, copies) << "rank " << rank; + EXPECT_EQ(runs[rank].bytes, bytes) << "rank " << rank; + EXPECT_EQ(runs[rank].slotEpochs, epochs) << "rank " << rank; + } + } + + // Copies in rounds of two or more ops; a one-op round takes the single-stream fallback even with copy streams. + static size_t MultiOpRoundCopies(const RankRun& run) { + std::map roundOps; + for (uint32_t epoch : run.slotEpochs) { + if (epoch != 0) { + ++roundOps[epoch]; + } + } + size_t copies = 0; + for (const auto& [epoch, ops] : roundOps) { + if (ops > 1) { + copies += ops; + } + } + return copies; + } + + // Guards the parameter itself: each arm must really reach the copy API it is named for. + void ExpectLaunchPathTaken(const std::vector& runs) { + for (size_t rank = 0; rank < runs.size(); ++rank) { + const size_t expectedBatch = GetParam() == LaunchPath::kBatchAsync ? runs[rank].copies : 0; + EXPECT_EQ(runs[rank].batchCopies, expectedBatch) << "rank " << rank; + const size_t expectedCopyStream = GetParam() == LaunchPath::kMultiStream ? MultiOpRoundCopies(runs[rank]) : 0; + EXPECT_EQ(runs[rank].copyStreamCopies, expectedCopyStream) << "rank " << rank; + } + } + + // Ready and complete barriers around the transfers, plus one per further round of freq destination slots. + static uint32_t RoundSyncBarriers(int nRanks, uint32_t freq) { return 2 + (nRanks - 1) / freq; } + + // Graph capture rounds up to an even count, so back-to-back replays keep alternating the ready/complete flags. + void ExpectRoundSyncBarriers(const std::vector& runs, int nRanks, uint32_t freq) { + uint32_t expected = RoundSyncBarriers(nRanks, freq); + if (GetParam() == LaunchPath::kGraphCapture && expected % 2 != 0) { + ++expected; + } + EXPECT_EQ(Barriers(runs), std::vector(nRanks, expected)); + } + + // hipStreamBatchMemOp calls per barrier: the wait batch, plus the flag write and the reset batch under graph capture. + int SyncCallsPerBarrier() const { return GetParam() == LaunchPath::kGraphCapture ? 3 : 1; } + + RankRun current_; + ncclComm* activeComm_ = nullptr; + size_t failCopyCall_ = 0; // 1-based copy API call on each rank that fails; 0 never fails + int failSyncCall_ = 0; // 1-based hipStreamBatchMemOp call on each rank that fails; 0 never fails +}; + +TEST_P(CeAlltoAllvSyncMicrotest, Sparse16Ranks_EqualTotalsAt512MiB_EveryRankIssuesSameBarrierCount) { + const int kNRanks = 16; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + const SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 6, 5), CE_COLL_INTRA_BATCH_SYNC_MSG_THRESHOLD); + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, Sparse32Ranks_EveryRankOverFreqButUnevenRounds_EveryRankIssuesSameBarrierCount) { + const int kNRanks = 32; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + const SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 12, 7), 12252240ull * 64); + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, Sparse4RanksFreq2_EqualTotalsOver512MiB_EveryRankIssuesSameBarrierCount) { + const int kNRanks = 4; + const uint32_t kFreq = 2; + const SizeMatrix send = EqualTotalsMatrix({2, 3, 4, 2}, 12 * 48 * kMiB); + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, OneHeavyRank_OthersUnderThreshold_EveryRankJoinsTheRoundBarriers) { + const int kNRanks = 16; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + for (int heavyRank : {0, kNRanks - 1}) { + SCOPED_TRACE("heavy rank " + std::to_string(heavyRank)); + SizeMatrix send = EqualTotalsMatrix(std::vector(kNRanks, 4), 4 * kMiB); + for (int j = 0; j < 10; ++j) { + send[heavyRank][(heavyRank + j) % kNRanks] = 64 * kMiB; + } + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); + } +} + +TEST_P(CeAlltoAllvSyncMicrotest, Sparse16Ranks_TotalsUnder512MiB_OnlyReadyAndCompleteBarriers) { + const int kNRanks = 16; + const SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 6, 5), 2520 * 100 * 1024); + + const std::vector runs = Run(send, CE_COLL_INTRA_BATCH_SYNC_FREQ); + + EXPECT_EQ(Barriers(runs), std::vector(kNRanks, 2)); + ExpectRowsCopied(send, runs, kNRanks); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, ManyPeersEachUnder512MiB_NoRankCrossesBothGates_OnlyReadyAndCompleteBarriers) { + const int kNRanks = 16; + const SizeMatrix send = EqualTotalsMatrix(std::vector(kNRanks, kNRanks), kNRanks * 31 * kMiB); + + const std::vector runs = Run(send, CE_COLL_INTRA_BATCH_SYNC_FREQ); + + EXPECT_EQ(Barriers(runs), std::vector(kNRanks, 2)); + ExpectRowsCopied(send, runs, kNRanks); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, Dense16Ranks_Over512MiB_EveryRankIssuesSameBarrierCount) { + const int kNRanks = 16; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + const SizeMatrix send = EqualTotalsMatrix(std::vector(kNRanks, kNRanks), kNRanks * 40 * kMiB); + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, SixteenRanks_EachSendsToExactlyFreqPeers_OnlyReadyAndCompleteBarriers) { + const int kNRanks = 16; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + const SizeMatrix send = EqualTotalsMatrix(std::vector(kNRanks, kFreq), 1024 * kMiB); + + const std::vector runs = Run(send, kFreq); + + EXPECT_EQ(Barriers(runs), std::vector(kNRanks, 2)); + ExpectRowsCopied(send, runs, kNRanks); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, SeventeenRanksFreq8_EvenRoundsPartialLastRound_EveryRankIssuesSameBarrierCount) { + const int kNRanks = 17; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + const SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 9, 9), CE_COLL_INTRA_BATCH_SYNC_MSG_THRESHOLD); + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, OneRankSendsNothing_OthersCrossBothGates_SilentRankIssuesSameBarrierCount) { + const int kNRanks = 16; + const int kSilentRank = 5; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 9, 5), CE_COLL_INTRA_BATCH_SYNC_MSG_THRESHOLD); + send[kSilentRank].assign(kNRanks, 0); + + const std::vector runs = Run(send, kFreq); + + ExpectRoundSyncBarriers(runs, kNRanks, kFreq); + ExpectRowsCopied(send, runs, kFreq); + ExpectLaunchPathTaken(runs); +} + +TEST_P(CeAlltoAllvSyncMicrotest, FirstRoundCopyFails_ReturnsTheErrorBeforeTheRoundBarrier) { + const int kNRanks = 16; + const SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 9, 5), CE_COLL_INTRA_BATCH_SYNC_MSG_THRESHOLD); + failCopyCall_ = 1; + + const std::vector runs = Run(send, CE_COLL_INTRA_BATCH_SYNC_FREQ, ncclUnhandledCudaError); + + EXPECT_EQ(Barriers(runs), std::vector(kNRanks, 1)); +} + +TEST_P(CeAlltoAllvSyncMicrotest, FirstRoundBarrierFails_ReturnsTheErrorWithoutLaunchingTheNextRound) { + const int kNRanks = 16; + const uint32_t kFreq = CE_COLL_INTRA_BATCH_SYNC_FREQ; + const SizeMatrix send = EqualTotalsMatrix(CyclingPeerCounts(kNRanks, 9, 5), CE_COLL_INTRA_BATCH_SYNC_MSG_THRESHOLD); + failSyncCall_ = SyncCallsPerBarrier() + 1; + + const std::vector runs = Run(send, kFreq, ncclUnhandledCudaError); + + EXPECT_EQ(Barriers(runs), std::vector(kNRanks, 2)); + for (int rank = 0; rank < kNRanks; ++rank) { + EXPECT_EQ(runs[rank].copies, kFreq) << "rank " << rank; + } +} + +INSTANTIATE_TEST_SUITE_P(LaunchPaths, CeAlltoAllvSyncMicrotest, + ::testing::Values(LaunchPath::kSingleStream, LaunchPath::kMultiStream, LaunchPath::kBatchAsync, + LaunchPath::kLegacyNullStream, LaunchPath::kGraphCapture), + LaunchPathName); + } // namespace diff --git a/projects/rccl/test/host/enqueue-test.cc b/projects/rccl/test/host/enqueue-test.cc index ef4dbca0a008..1e529954b5a2 100644 --- a/projects/rccl/test/host/enqueue-test.cc +++ b/projects/rccl/test/host/enqueue-test.cc @@ -3506,6 +3506,10 @@ TEST_F(EnqueueMicrotest, TopoGetAlgoInfo_ConsultsEveryRcclTuningHookExactlyOnce) // a lie: deleting either call site was invisible. EXPECT_EQ(1, g_rcclOverrideAlgorithmCalls) << "rcclOverrideAlgorithm (:2733)"; EXPECT_EQ(1, g_rcclOverrideProtocolCalls) << "rcclOverrideProtocol (:2734)"; + EXPECT_EQ(1, g_rcclUpdateCollectiveAlgorithmCalls) << "rcclUpdateCollectiveAlgorithm"; + EXPECT_TRUE(g_rcclUpdateCollectiveAlgorithmBeforeProtocol) + << "rcclUpdateCollectiveAlgorithm must run before rcclUpdateCollectiveProtocol picks the Tree protocol"; + EXPECT_EQ(1, g_rcclUpdateCollectiveChannelsCalls) << "rcclUpdateCollectiveChannels"; } // LATENT BUG (enqueue.cc:2659): rcclOverrideChannels returns ncclResult_t, but diff --git a/projects/rccl/test/host/fakes/dev_runtime_fakes.cc b/projects/rccl/test/host/fakes/dev_runtime_fakes.cc index 89ba710e787e..a7149af3a2e6 100644 --- a/projects/rccl/test/host/fakes/dev_runtime_fakes.cc +++ b/projects/rccl/test/host/fakes/dev_runtime_fakes.cc @@ -17,6 +17,7 @@ ASSERT_HOOK_MATCHES_PROD(g_devrFindWindow, ncclDevrFindWindow); ASSERT_HOOK_MATCHES_PROD(g_devrWindowIsMultiSegment, ncclDevrWindowIsMultiSegment); ASSERT_HOOK_MATCHES_PROD(g_devrWindowHasSysmemSegment, ncclDevrWindowHasSysmemSegment); ASSERT_HOOK_MATCHES_PROD(g_devrInitOnce, ncclDevrInitOnce); +ASSERT_HOOK_MATCHES_PROD(g_devrGetLsaRankPtr, ncclDevrGetLsaRankPtr); ASSERT_HOOK_MATCHES_PROD(g_devrWindowRegisterInGroup, ncclDevrWindowRegisterInGroup); ASSERT_HOOK_MATCHES_PROD(g_devrNcclCommWindowDeregister, ncclCommWindowDeregister); #undef ASSERT_HOOK_MATCHES_PROD @@ -51,6 +52,16 @@ static ncclResult_t DefaultDevrInitOnce(struct ncclComm*) { return ncclSuccess; std::function g_devrInitOnce = DefaultDevrInitOnce; ncclResult_t ncclDevrInitOnce(struct ncclComm* comm) { return g_devrInitOnce(comm); } +static ncclResult_t DefaultDevrGetLsaRankPtr(struct ncclComm*, struct ncclDevrWindow*, size_t, int, void**) { + FailLoudUnfaked("dev_runtime_fakes", "ncclDevrGetLsaRankPtr"); +} +std::function g_devrGetLsaRankPtr = + DefaultDevrGetLsaRankPtr; +ncclResult_t ncclDevrGetLsaRankPtr(struct ncclComm* comm, struct ncclDevrWindow* winHost, size_t offset, int lsaRank, + void** outPtr) { + return g_devrGetLsaRankPtr(comm, winHost, offset, lsaRank, outPtr); +} + static ncclResult_t DefaultDevrWindowRegisterInGroup(struct ncclComm*, void*, size_t, int, ncclWindow_t*) { FailLoudUnfaked("dev_runtime_fakes", "ncclDevrWindowRegisterInGroup"); } @@ -70,9 +81,6 @@ ncclResult_t ncclCommWindowDeregister(ncclComm_t comm, ncclWindow_t win) { } // Floors for the LSA addressing ce_coll.cc's copy paths reach. -ncclResult_t ncclDevrGetLsaRankPtr(struct ncclComm*, struct ncclDevrWindow*, size_t, int, void**) { - FailLoudUnfaked("dev_runtime_fakes", "ncclDevrGetLsaRankPtr"); -} ncclResult_t ncclDevrWorldToLsaRank(struct ncclComm*, int, int*) { FailLoudUnfaked("dev_runtime_fakes", "ncclDevrWorldToLsaRank"); } @@ -87,6 +95,7 @@ void ResetDevRuntimeFakes() { g_devrWindowIsMultiSegment = DefaultDevrWindowIsMultiSegment; g_devrWindowHasSysmemSegment = DefaultDevrWindowHasSysmemSegment; g_devrInitOnce = DefaultDevrInitOnce; + g_devrGetLsaRankPtr = DefaultDevrGetLsaRankPtr; g_devrWindowRegisterInGroup = DefaultDevrWindowRegisterInGroup; g_devrNcclCommWindowDeregister = DefaultCommWindowDeregister; } diff --git a/projects/rccl/test/host/fakes/dev_runtime_fakes.h b/projects/rccl/test/host/fakes/dev_runtime_fakes.h index fd98e42849c7..f4851ef81e4c 100644 --- a/projects/rccl/test/host/fakes/dev_runtime_fakes.h +++ b/projects/rccl/test/host/fakes/dev_runtime_fakes.h @@ -11,6 +11,7 @@ #ifndef RCCL_TEST_HOST_DEV_RUNTIME_FAKES_H_ #define RCCL_TEST_HOST_DEV_RUNTIME_FAKES_H_ +#include #include #include "nccl.h" @@ -28,6 +29,10 @@ extern bool g_devrWindowHasSysmemSegmentValue; // never sets symmetricSupport with a null peerInfo (dev_runtime.cc's own real error condition). extern std::function g_devrInitOnce; +// ncclDevrGetLsaRankPtr: fail-loud by default, since a peer address only exists for a window layout the test defines. +extern std::function + g_devrGetLsaRankPtr; // Fail-loud default: a window registration is collective, so a test must opt in to reaching one. extern std::function g_devrWindowRegisterInGroup; // Fail-loud default; same spelling as dev_runtime_micro_fakes.h, which no target links alongside this file. diff --git a/projects/rccl/test/host/fakes/hip_fakes.cc b/projects/rccl/test/host/fakes/hip_fakes.cc index 06258373ebfe..b2eac23f048b 100644 --- a/projects/rccl/test/host/fakes/hip_fakes.cc +++ b/projects/rccl/test/host/fakes/hip_fakes.cc @@ -55,6 +55,9 @@ ASSERT_HOOK_MATCHES_PROD(g_hipEventRecord, hipEventRecord); ASSERT_HOOK_MATCHES_PROD(g_hipStreamBatchMemOp, hipStreamBatchMemOp); ASSERT_HOOK_MATCHES_PROD(g_hipStreamWriteValue64, hipStreamWriteValue64); ASSERT_HOOK_MATCHES_PROD(g_hipStreamWaitValue64, hipStreamWaitValue64); +#ifdef CE_BATCH_ASYNC_SUPPORTED +ASSERT_HOOK_MATCHES_PROD(g_hipMemcpyBatchAsync, hipMemcpyBatchAsync); +#endif #undef ASSERT_HOOK_MATCHES_PROD @@ -546,6 +549,17 @@ static hipError_t DefaultHipStreamBatchMemOp(hipStream_t, unsigned int, std::function g_hipStreamBatchMemOp = DefaultHipStreamBatchMemOp; +#ifdef CE_BATCH_ASYNC_SUPPORTED +static hipError_t DefaultHipMemcpyBatchAsync(void**, void**, size_t*, size_t, hipMemcpyAttributes*, size_t*, + size_t, size_t*, hipStream_t) +{ + return g_hipAsyncOpsResult; +} +std::function + g_hipMemcpyBatchAsync = DefaultHipMemcpyBatchAsync; +#endif + static hipError_t DefaultHipStreamWriteValue64(hipStream_t, void*, std::uint64_t, unsigned int) { @@ -627,6 +641,9 @@ void ResetHipFakes() g_hipEventQueryRequiresRecord = false; g_recordedEvents.clear(); g_hipStreamBatchMemOp = DefaultHipStreamBatchMemOp; +#ifdef CE_BATCH_ASYNC_SUPPORTED + g_hipMemcpyBatchAsync = DefaultHipMemcpyBatchAsync; +#endif g_hipStreamWriteValue64 = DefaultHipStreamWriteValue64; g_hipStreamWaitValue64 = DefaultHipStreamWaitValue64; } @@ -927,10 +944,10 @@ hipError_t hipStreamWriteValue32(hipStream_t, void*, std::uint32_t, unsigned int // ce_coll.cc's batch path; in-RCCL builds inherit CE_BATCH_ASYNC_SUPPORTED, standalone ones compile the call out. #ifdef CE_BATCH_ASYNC_SUPPORTED -hipError_t hipMemcpyBatchAsync(void**, void**, size_t*, size_t, hipMemcpyAttributes*, size_t*, size_t, size_t*, - hipStream_t) +hipError_t hipMemcpyBatchAsync(void** dsts, void** srcs, size_t* sizes, size_t count, hipMemcpyAttributes* attrs, + size_t* attrsIdxs, size_t numAttrs, size_t* failIdx, hipStream_t stream) { - FailLoudUnfaked("hip_fakes", "hipMemcpyBatchAsync"); + return g_hipMemcpyBatchAsync(dsts, srcs, sizes, count, attrs, attrsIdxs, numAttrs, failIdx, stream); } #endif diff --git a/projects/rccl/test/host/fakes/hip_fakes.h b/projects/rccl/test/host/fakes/hip_fakes.h index 4e8bfbac5026..0a8f107d5025 100644 --- a/projects/rccl/test/host/fakes/hip_fakes.h +++ b/projects/rccl/test/host/fakes/hip_fakes.h @@ -217,6 +217,12 @@ extern std::function g_hipStreamWaitValue64; +#ifdef CE_BATCH_ASYNC_SUPPORTED +extern std::function + g_hipMemcpyBatchAsync; +#endif // Restore the HIP controllable seams above to their defaults. Called by // ResetP2pFakes(); exposed for tests that only touch HIP hooks. diff --git a/projects/rccl/test/host/fakes/init_fakes.cc b/projects/rccl/test/host/fakes/init_fakes.cc index 3ae52f83ca50..d3f2a56f50ad 100644 --- a/projects/rccl/test/host/fakes/init_fakes.cc +++ b/projects/rccl/test/host/fakes/init_fakes.cc @@ -98,6 +98,7 @@ void ResetInitFakes() { ResetNcclFakes(); ResetNcclStubs(); ResetOsFakes(); + ResetRcclArchThresholdsFakes(); ResetRcclWrapFakes(); ResetRecorderFakes(); ResetRocmWrapFakes(); diff --git a/projects/rccl/test/host/fakes/init_fakes.h b/projects/rccl/test/host/fakes/init_fakes.h index d1fabb720f51..d004ee3aa190 100644 --- a/projects/rccl/test/host/fakes/init_fakes.h +++ b/projects/rccl/test/host/fakes/init_fakes.h @@ -24,6 +24,7 @@ #include "nccl_fakes.h" #include "nccl_stubs.h" // core/lifecycle stubs + data symbols #include "os_fakes.h" // src/os/linux.cc +#include "rccl_arch_thresholds_fakes.h" // src/rccl_arch_thresholds.cc #include "rccl_wrap_fakes.h" // src/rccl_wrap.cc #include "recorder_fakes.h" // src/recorder.cc #include "rocmwrap_fakes.h" // src/misc/rocmwrap.cc diff --git a/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.cc b/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.cc index e994d2ffaac5..13222036c04f 100644 --- a/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.cc +++ b/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.cc @@ -5,13 +5,25 @@ ************************************************************************/ // Fake for src/rccl_arch_thresholds.cc. -// Host-only test binaries do not link the real TU (it pulls in net.h / archinfo -// symbols that are already provided by other stubs), so this stub satisfies the -// linker. Returning nullptr is the correct "no arch table" sentinel -- every -// caller in production null-checks before dereferencing. +// Host-only test binaries do not link the real TU, so this stub satisfies the linker. +// By default it returns nullptr, the "no arch table" sentinel every production caller +// null-checks; a test can install a table per partition to see which one a comm got. -#include "rccl_arch_thresholds.h" +#include "rccl_arch_thresholds_fakes.h" -const rcclArchThresholds* rcclGetArchThresholds(const char* /*gcn*/) { - return nullptr; +int g_archThresholdsCalls = 0; +rcclTuningPartition g_lastArchThresholdsPartition = RCCL_TUNING_PARTITION_SPX; +const rcclArchThresholds* g_archThresholdsTable[2] = {nullptr, nullptr}; + +const rcclArchThresholds* rcclGetArchThresholds(const char* /*gcn*/, rcclTuningPartition partition) { + ++g_archThresholdsCalls; + g_lastArchThresholdsPartition = partition; + return g_archThresholdsTable[partition]; +} + +void ResetRcclArchThresholdsFakes() { + g_archThresholdsCalls = 0; + g_lastArchThresholdsPartition = RCCL_TUNING_PARTITION_SPX; + g_archThresholdsTable[RCCL_TUNING_PARTITION_SPX] = nullptr; + g_archThresholdsTable[RCCL_TUNING_PARTITION_DPX] = nullptr; } diff --git a/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.h b/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.h new file mode 100644 index 000000000000..d7d4529bfb49 --- /dev/null +++ b/projects/rccl/test/host/fakes/rccl_arch_thresholds_fakes.h @@ -0,0 +1,23 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Seams for the src/rccl_arch_thresholds.cc fake. + +#ifndef RCCL_TEST_HOST_RCCL_ARCH_THRESHOLDS_FAKES_H_ +#define RCCL_TEST_HOST_RCCL_ARCH_THRESHOLDS_FAKES_H_ + +#include "rccl_arch_thresholds.h" + +// Calls to rcclGetArchThresholds and the partition the last one asked for. +extern int g_archThresholdsCalls; +extern rcclTuningPartition g_lastArchThresholdsPartition; +// What rcclGetArchThresholds returns for each rcclTuningPartition (nullptr until a test sets it), +// so a test can check which table ended up in comm->archThresholds. +extern const rcclArchThresholds* g_archThresholdsTable[2]; + +void ResetRcclArchThresholdsFakes(); + +#endif // RCCL_TEST_HOST_RCCL_ARCH_THRESHOLDS_FAKES_H_ diff --git a/projects/rccl/test/host/fakes/rccl_wrap_fakes.cc b/projects/rccl/test/host/fakes/rccl_wrap_fakes.cc index b01a153ec881..7e97dcfd6438 100644 --- a/projects/rccl/test/host/fakes/rccl_wrap_fakes.cc +++ b/projects/rccl/test/host/fakes/rccl_wrap_fakes.cc @@ -40,6 +40,17 @@ void rcclSetPipelining(struct ncclComm* comm, size_t const& nBytes, struct ncclT ++g_rcclSetPipeliningCalls; g_rcclSetPipelining(comm, nBytes, info); } +int g_rcclUpdateCollectiveAlgorithmCalls = 0; +bool g_rcclUpdateCollectiveAlgorithmBeforeProtocol = false; +int g_rcclUpdateCollectiveChannelsCalls = 0; +void rcclUpdateCollectiveAlgorithm(struct ncclComm*, size_t, struct ncclTaskColl*, float (*)[NCCL_NUM_PROTOCOLS]) { + ++g_rcclUpdateCollectiveAlgorithmCalls; + // Production requires the algorithm hook to run before the protocol hook. + g_rcclUpdateCollectiveAlgorithmBeforeProtocol = g_rcclUpdateCollectiveProtocolCalls == 0; +} +void rcclUpdateCollectiveChannels(struct ncclComm*, size_t, const struct ncclTaskColl*, int&) { + ++g_rcclUpdateCollectiveChannelsCalls; +} void rcclUpdateThreadThreshold(struct ncclComm*, size_t const&, struct ncclTaskColl*, int&) { ++g_rcclUpdateThreadThresholdCalls; } @@ -164,6 +175,9 @@ void ResetRcclWrapFakes() { g_rcclParamCeAllReduce = -1; g_rcclParamCeArStagingBytes = -1; g_rcclUpdateCollectiveProtocolCalls = 0; + g_rcclUpdateCollectiveAlgorithmCalls = 0; + g_rcclUpdateCollectiveAlgorithmBeforeProtocol = false; + g_rcclUpdateCollectiveChannelsCalls = 0; g_rcclSetPipeliningCalls = 0; g_rcclUpdateThreadThresholdCalls = 0; g_rcclOptThreadBlockSizeCalls = 0; diff --git a/projects/rccl/test/host/fakes/rccl_wrap_fakes.h b/projects/rccl/test/host/fakes/rccl_wrap_fakes.h index a9f9e2dc9830..b3a27c911829 100644 --- a/projects/rccl/test/host/fakes/rccl_wrap_fakes.h +++ b/projects/rccl/test/host/fakes/rccl_wrap_fakes.h @@ -40,6 +40,9 @@ extern int64_t g_rcclParamDirectReduceScatterThreshold; // UNDRIVEN // Call counters for the no-op tuning hooks: a no-op that was never called and one // that was look identical without these, so a dropped call site would be silent. extern int g_rcclUpdateCollectiveProtocolCalls; +extern int g_rcclUpdateCollectiveAlgorithmCalls; +extern bool g_rcclUpdateCollectiveAlgorithmBeforeProtocol; +extern int g_rcclUpdateCollectiveChannelsCalls; extern int g_rcclSetPipeliningCalls; extern int g_rcclUpdateThreadThresholdCalls; extern int g_rcclOptThreadBlockSizeCalls; diff --git a/projects/rccl/test/host/init-test.cc b/projects/rccl/test/host/init-test.cc index 4da70f57a9cb..355669285bbb 100644 --- a/projects/rccl/test/host/init-test.cc +++ b/projects/rccl/test/host/init-test.cc @@ -1818,11 +1818,14 @@ TEST_F(InitMicrotest, FillInfo_FunctionAboveMloPartMax_LeavesMloPartUndefined) { FillInfoComm c; c.get()->busId = kHighFnBusId; g_pciDeviceClass = ""; // no sysfs entry, so only the fn bound can reject the stamp + g_pciComputePartition = "DPX"; ncclPeerInfo info{}; EXPECT_EQ(ncclSuccess, fillInfo(c.get(), &info, 0)); // 0xf does not fit the 3-bit overlay index, so it must not be stamped even though sysfs is empty. EXPECT_EQ(NCCL_TOPO_UNDEF, info.mloPart); EXPECT_EQ(0, g_pciDeviceClassCalls); + // The tuning request does not depend on the MLOPart bound: the mode is still read and used. + EXPECT_EQ(RCCL_TUNING_PARTITION_DPX, info.tuningPartition); } TEST_F(InitMicrotest, FillInfo_CpxPartitionZero_StampsMloPartZero) { @@ -1884,6 +1887,179 @@ TEST_F(InitMicrotest, FillInfo_NoPartitionModeReported_FunctionZeroLeavesMloPart EXPECT_EQ(NCCL_TOPO_UNDEF, info.mloPart); } +// RCCL_TUNING_PARTITION_MODE: each rank asks for the gfx1250 tuning table of the mode its GPU +// reports, from the same sysfs read the MLOPart probe uses; only DPX asks for the DPX table. The +// env var forces a table either way, and "auto" or an empty value follow the GPU. +TEST_F(InitMicrotest, FillInfo_TuningPartitionRequest) { + struct Case { + const char* env; // nullptr = unset + const char* gpu; + int want; + }; + const Case cases[] = {{nullptr, "SPX", RCCL_TUNING_PARTITION_SPX}, + {nullptr, "DPX", RCCL_TUNING_PARTITION_DPX}, + {nullptr, "CPX", RCCL_TUNING_PARTITION_SPX}, + {nullptr, "", RCCL_TUNING_PARTITION_SPX}, + {"SPX", "DPX", RCCL_TUNING_PARTITION_SPX}, + {"DPX", "SPX", RCCL_TUNING_PARTITION_DPX}, + {"dpx", "", RCCL_TUNING_PARTITION_DPX}, + {"Auto", "DPX", RCCL_TUNING_PARTITION_DPX}, + {"", "DPX", RCCL_TUNING_PARTITION_DPX}}; + for (const Case& tc : cases) { + const std::string label = std::string("env '") + (tc.env ? tc.env : "") + "', GPU '" + tc.gpu + "'"; + SetMicroEnv("RCCL_TUNING_PARTITION_MODE", tc.env); + FillInfoComm c; + c.get()->busId = kAliasBusId; + g_pciComputePartition = tc.gpu; + g_pciComputePartitionCalls = 0; + ncclPeerInfo info{}; + info.tuningPartition = -1; + ncclResult_t res = ncclInternalError; + ScopedDebugLogging dbg(NCCL_LOG_INFO, NCCL_ALL); + const std::string log = RcclUnitTesting::CaptureLog([&] { res = fillInfo(c.get(), &info, 0); }); + EXPECT_EQ(ncclSuccess, res); + EXPECT_EQ(tc.want, info.tuningPartition) << label; + EXPECT_EQ(1, g_pciComputePartitionCalls) << "one read serves MLOPart and tuning, " << label; + EXPECT_EQ("0000:11:00.0", g_lastPciComputePartitionBusId) << "the mode lives on function 0"; + EXPECT_FALSE(LogHas(log, "RCCL_TUNING_PARTITION_MODE")) << "a valid value must not be reported:\n" << log; + } +} + +// An unknown value is ignored with one notice per process, and the GPU's mode decides. Isolated: +// the report-once flag is a function-local static. +TEST_F(InitMicrotestIsolated, FillInfo_InvalidTuningPartitionMode_ReportsOnceAndFollowsGpuMode) { + RUN_ISOLATED_TEST( + "Init_FillInfo_InvalidTuningPartitionMode", + []() { + SetMicroEnv("RCCL_TUNING_PARTITION_MODE", "QPX"); + g_pciComputePartition = "DPX"; + FillInfoComm c; + c.get()->busId = kPhysGpuBusId; + ncclPeerInfo first{}; + ncclPeerInfo second{}; + ncclResult_t r1 = ncclInternalError; + ncclResult_t r2 = ncclInternalError; + ScopedDebugLogging dbg(NCCL_LOG_INFO, NCCL_ALL); + const std::string log = RcclUnitTesting::CaptureLog([&] { + r1 = fillInfo(c.get(), &first, 0); + r2 = fillInfo(c.get(), &second, 0); + }); + EXPECT_EQ(ncclSuccess, r1); + EXPECT_EQ(ncclSuccess, r2); + EXPECT_EQ(RCCL_TUNING_PARTITION_DPX, first.tuningPartition); + EXPECT_EQ(RCCL_TUNING_PARTITION_DPX, second.tuningPartition); + const std::string needle = "Invalid value QPX for RCCL_TUNING_PARTITION_MODE"; + const size_t at = log.find(needle); + EXPECT_NE(std::string::npos, at) << "actual log:\n" << log; + EXPECT_EQ(std::string::npos, log.find(needle, at + needle.size())) << "reported twice:\n" << log; + }); +} + +namespace { +// AllGather1 output for ranks that asked for `requests`. +std::vector PeersAsking(std::initializer_list requests) { + std::vector peers(requests.size()); + size_t i = 0; + for (int request : requests) peers[i++].tuningPartition = request; + return peers; +} + +rcclTuningPartition TuningPartitionFor(std::initializer_list requests, bool* mixed) { + std::vector peers = PeersAsking(requests); + return rcclTuningPartitionForComm(peers.data(), static_cast(peers.size()), mixed); +} + +// Distinct tables for the rcclGetArchThresholds fake to hand out, so a test sees which one a comm got. +const rcclArchThresholds kFakeSpxTable{}; +const rcclArchThresholds kFakeDpxTable{}; +void InstallFakeTuningTables() { + g_archThresholdsTable[RCCL_TUNING_PARTITION_SPX] = &kFakeSpxTable; + g_archThresholdsTable[RCCL_TUNING_PARTITION_DPX] = &kFakeDpxTable; +} +} // namespace + +// Every rank must read the same table, so the comm is DPX only when all ranks ask for DPX. +TEST_F(InitMicrotest, TuningPartitionForComm_DpxOnlyWhenEveryRankAsks) { + constexpr int kSpx = RCCL_TUNING_PARTITION_SPX; + constexpr int kDpx = RCCL_TUNING_PARTITION_DPX; + bool mixed = true; + EXPECT_EQ(RCCL_TUNING_PARTITION_DPX, TuningPartitionFor({kDpx, kDpx, kDpx}, &mixed)); + EXPECT_FALSE(mixed); + EXPECT_EQ(RCCL_TUNING_PARTITION_DPX, TuningPartitionFor({kDpx}, &mixed)); + EXPECT_FALSE(mixed); + EXPECT_EQ(RCCL_TUNING_PARTITION_SPX, TuningPartitionFor({kSpx, kSpx}, &mixed)); + EXPECT_FALSE(mixed); + EXPECT_EQ(RCCL_TUNING_PARTITION_SPX, TuningPartitionFor({kDpx, kSpx, kDpx}, &mixed)); + EXPECT_TRUE(mixed); + EXPECT_EQ(RCCL_TUNING_PARTITION_SPX, TuningPartitionFor({kSpx, kDpx}, &mixed)); + EXPECT_TRUE(mixed); +} + +namespace { +// A gfx1250 comm whose ranks asked for `requests`, as rcclSetTuningPartition sees it after AllGather1. +struct TuningPartitionComm { + explicit TuningPartitionComm(std::initializer_list requests, const char* arch = "gfx1250", int rank = 0) + : peers(PeersAsking(requests)) { + comm.get()->peerInfo = peers.data(); + comm.get()->nRanks = static_cast(peers.size()); + comm.get()->rank = rank; + comm.get()->archName = const_cast(arch); + comm.get()->archThresholds = nullptr; + InstallFakeTuningTables(); + } + ~TuningPartitionComm() { + comm.get()->peerInfo = nullptr; + comm.get()->archName = nullptr; + comm.get()->archThresholds = nullptr; + } + std::string Run() { + ScopedDebugLogging dbg(NCCL_LOG_INFO, NCCL_ALL); + return RcclUnitTesting::CaptureLog([&] { rcclSetTuningPartition(comm.get()); }); + } + FillInfoComm comm; + std::vector peers; +}; +} // namespace + +TEST_F(InitMicrotest, SetTuningPartition_AllDpxPicksTheDpxTable) { + TuningPartitionComm t({RCCL_TUNING_PARTITION_DPX, RCCL_TUNING_PARTITION_DPX}); + const std::string log = t.Run(); + EXPECT_EQ(&kFakeDpxTable, t.comm.get()->archThresholds); + EXPECT_EQ(1, g_archThresholdsCalls); + EXPECT_TRUE(LogHas(log, "gfx1250 tuning table: DPX")) << "actual log:\n" << log; + EXPECT_FALSE(LogHas(log, "ask for different tables")) << "actual log:\n" << log; +} + +TEST_F(InitMicrotest, SetTuningPartition_MixedRanksUseSpxAndSayWhy) { + TuningPartitionComm t({RCCL_TUNING_PARTITION_DPX, RCCL_TUNING_PARTITION_SPX}); + const std::string log = t.Run(); + EXPECT_EQ(&kFakeSpxTable, t.comm.get()->archThresholds); + EXPECT_TRUE(LogHas(log, "ask for different tables (RCCL_TUNING_PARTITION_MODE")) << "actual log:\n" << log; + EXPECT_TRUE(LogHas(log, "gfx1250 tuning table: SPX")) << "actual log:\n" << log; +} + +// Only rank 0 of a gfx1250 comm logs, but every rank and arch still re-picks its table. +TEST_F(InitMicrotest, SetTuningPartition_OnlyRankZeroOfGfx1250Logs) { + { + TuningPartitionComm t({RCCL_TUNING_PARTITION_DPX, RCCL_TUNING_PARTITION_DPX}, "gfx1250", /*rank=*/1); + const std::string log = t.Run(); + EXPECT_EQ(&kFakeDpxTable, t.comm.get()->archThresholds); + EXPECT_FALSE(LogHas(log, "tuning table")) << "actual log:\n" << log; + } + { + TuningPartitionComm t({RCCL_TUNING_PARTITION_DPX, RCCL_TUNING_PARTITION_SPX}, "gfx1250", /*rank=*/1); + const std::string log = t.Run(); + EXPECT_EQ(&kFakeSpxTable, t.comm.get()->archThresholds); + EXPECT_FALSE(LogHas(log, "tuning table")) << "a mixed comm is reported by rank 0 only:\n" << log; + } + { + TuningPartitionComm t({RCCL_TUNING_PARTITION_DPX, RCCL_TUNING_PARTITION_SPX}, "gfx950"); + const std::string log = t.Run(); + EXPECT_EQ(&kFakeSpxTable, t.comm.get()->archThresholds); + EXPECT_FALSE(LogHas(log, "tuning")) << "actual log:\n" << log; + } +} + TEST_F(InitMicrotest, CommInitAll_GetDeviceFault_StopsBeforeCount) { g_hipGetDevice = [](int*) { return hipErrorInvalidValue; }; g_hipGetDeviceCount = [](int*) -> hipError_t { @@ -2533,6 +2709,7 @@ struct PeerSpec { int uuidTag = -1; // -1 = its own UUID; equal values (>=1) make two peers the same GPU int64_t busId = 0; // equal values are the same PCI function, i.e. the same partition int compCap = -1; // -1 = same as self + int tuningPartition = -1; // -1 = same as self ncclComm* comm = nullptr; // only read for same-process peers (:1527, :1549) }; @@ -2576,6 +2753,7 @@ void InstallPeerInfoAllGather(TransportsRankComm& c, std::vector specs reinterpret_cast(&info[i].gpuUuid)[0] = static_cast(s.uuidTag < 0 ? i + 1 : s.uuidTag); // 1-based: never matches self's zeros info[i].cudaCompCap = s.compCap < 0 ? self.cudaCompCap : s.compCap; + info[i].tuningPartition = s.tuningPartition < 0 ? self.tuningPartition : s.tuningPartition; info[i].comm = s.comm; } return ncclSuccess; @@ -2670,6 +2848,32 @@ TEST_F(InitMicrotest, InitTransportsRank_PeerWithMloPart_SetsHasMloPart) { EXPECT_TRUE(c.get()->hasMloPart); } +// RCCL_TUNING_PARTITION_MODE: after AllGather1 the comm re-picks its tuning table from every +// rank's request -- DPX only when all of them ask for it. +TEST_F(InitMicrotest, InitTransportsRank_EveryRankAsksDpx_PicksTheDpxTable) { + TransportsRankComm c(/*nRanks=*/4, /*rank=*/0); + g_pciComputePartition = "DPX"; // fillInfo's own slot asks for DPX; peers copy it + InstallPeerInfoAllGather(c, std::vector(4)); + InstallFakeTuningTables(); + c.get()->archThresholds = nullptr; + EXPECT_EQ(ncclRemoteError, initTransportsRank(c.get(), nullptr, c.timers())); + EXPECT_EQ(&kFakeDpxTable, c.get()->archThresholds); + EXPECT_EQ(1, g_archThresholdsCalls); +} + +// Rank 0 itself asks for DPX, so only counting every peer's slot can pick SPX. +TEST_F(InitMicrotest, InitTransportsRank_OnePeerAsksSpx_PicksTheSpxTable) { + TransportsRankComm c(/*nRanks=*/4, /*rank=*/0); + g_pciComputePartition = "DPX"; + std::vector specs(4); + specs[2].tuningPartition = RCCL_TUNING_PARTITION_SPX; + InstallPeerInfoAllGather(c, specs); + InstallFakeTuningTables(); + c.get()->archThresholds = &kFakeDpxTable; + EXPECT_EQ(ncclRemoteError, initTransportsRank(c.get(), nullptr, c.timers())); + EXPECT_EQ(&kFakeSpxTable, c.get()->archThresholds); +} + TEST_F(InitMicrotest, InitTransportsRank_NoPeerWithMloPart_LeavesHasMloPartUnset) { TransportsRankComm c(/*nRanks=*/4, /*rank=*/0); InstallPeerInfoAllGather(c, std::vector(4)); diff --git a/projects/rccl/test/host/wrap-test.cc b/projects/rccl/test/host/wrap-test.cc index 5bd1f8c092a8..284c2e8d4aff 100644 --- a/projects/rccl/test/host/wrap-test.cc +++ b/projects/rccl/test/host/wrap-test.cc @@ -1276,6 +1276,270 @@ TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_MultiNodeUsesTunedLLRange) }); } +namespace { +constexpr size_t kKiB = 1024, kMiB = 1024 * 1024; + +// The gfx1250 arch-table entries this change reads (rccl_arch_thresholds.cc). +// rccl-UnitTestsMicro links no real tables (the fake rcclGetArchThresholds returns +// null), so tests hand this copy to the comm; RcclDdaTierThresholds.* (GPU) checks +// that the real table resolves to the same caps. +rcclArchThresholds MakeRankScaledTable() { + rcclArchThresholds t{}; + t.ddaLLMax[ncclFuncAllReduce] = 32 * kMiB; + t.ddaLLMax[ncclFuncAllGather] = 128 * kKiB; + t.ddaLLMax[ncclFuncReduceScatter] = 16 * kMiB; + t.ddaLL128Max[ncclFuncAllReduce] = 256 * kMiB; + t.ddaLL128Max[ncclFuncAllGather] = 128 * kMiB; + t.ddaLL128Max[ncclFuncReduceScatter] = 2 * kMiB; + t.ddaVmmMax[ncclFuncAllReduce] = 16 * kMiB; + t.ddaLLMaxPerRank[ncclFuncAllReduce] = 1 * kMiB; + t.ddaLLMaxPerRank[ncclFuncAllGather] = 16 * kKiB; + t.ddaLLMaxPerRank[ncclFuncReduceScatter] = 256 * kKiB; + t.ddaLL128MaxPerRank[ncclFuncAllGather] = 16 * kMiB; + t.ddaLL128MaxPerRank[ncclFuncReduceScatter] = 512 * kKiB; + t.ddaVmmMaxPerRank[ncclFuncReduceScatter] = 2 * kMiB; + t.rankScaledMinRanks = 8; + t.rankScaledMaxBytes = 256 * kMiB; + t.ddaArVmmFirstMinRanks = 32; + t.ddaArVmmFirstMinBytes = 2 * kMiB; + t.ddaArVmmFirstMaxPerRank = 512 * kKiB; + t.ddaVmmBlockBudget = 2048; + t.ddaVmmMinBlocks = 32; + t.ddaVmmMaxBlocks = 128; + t.arRingLL128MaxPerRank = 32 * kMiB; + t.arTreeLL128Unlimited = true; + t.arTreeMaxPerRank = 1 * kMiB; + t.arTreeLLMaxPerRank = 32 * kKiB; + t.agRingLL128MaxPerRank = 64 * kMiB; + t.rsRingLL128MinPerRank = 512 * kKiB; + t.rsRingLL128MaxPerRank = 8 * kMiB; + t.rsRingLL128MaxPerRank1B = 32 * kMiB; + t.rsRingSimpleCappedMaxPerRank = 32 * kMiB; + t.rsRingSimpleCappedChannels = 128; + return t; +} + +const rcclArchThresholds* Gfx1250Table() { + static const rcclArchThresholds table = MakeRankScaledTable(); + return &table; +} + +// gfx1250 single-node comm carrying the arch table and tuning_model_11's AllReduce, +// AllGather and ReduceScatter LL / LL128 ranges. +constexpr uint64_t kGfx1250ArLL128Max = 16777216; +ncclComm* MakeGfx1250ProtocolComm(int nRanks) { + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->nNodes = 1; + comm->nRanks = nRanks; + comm->archThresholds = Gfx1250Table(); + comm->topo->ll128Enabled = true; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL][RCCL_PROTOCOL_MIN_IDX] = 0; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL][RCCL_PROTOCOL_MAX_IDX] = 131072; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MIN_IDX] = 131072; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MAX_IDX] = kGfx1250ArLL128Max; + for (int t : {RCCL_AG_TUNABLE, RCCL_RS_TUNABLE}) { + comm->minMaxLLRange[t][NCCL_PROTO_LL][RCCL_PROTOCOL_MIN_IDX] = 0; + comm->minMaxLLRange[t][NCCL_PROTO_LL][RCCL_PROTOCOL_MAX_IDX] = 1048576; + comm->minMaxLLRange[t][NCCL_PROTO_LL128][RCCL_PROTOCOL_MIN_IDX] = 1048576; + } + comm->minMaxLLRange[RCCL_AG_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MAX_IDX] = 16777216; + comm->minMaxLLRange[RCCL_RS_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MAX_IDX] = 4194304; + return comm; +} + +// Protocol rcclUpdateCollectiveProtocol picks for `func` / `algorithm`. AllReduce +// passes the whole message; AllGather and ReduceScatter pass the per-rank share +// (nBytes is nRanks x that), matching rcclGetSizePerRank(). +int ProtocolAfterUpdate(ncclComm* comm, ncclFunc_t func, int algorithm, size_t bytes, + ncclDataType_t dt = ncclFloat32) { + ncclTaskColl info{}; + info.func = func; + info.algorithm = algorithm; + info.protocol = NCCL_PROTO_SIMPLE; + info.datatype = dt; + const size_t nBytes = func == ncclFuncAllReduce ? bytes : bytes * (size_t)comm->nRanks; + rcclUpdateCollectiveProtocol(comm, nBytes, &info); + return info.protocol; +} +int ArProtocol(ncclComm* comm, int algorithm, size_t nBytes) { + return ProtocolAfterUpdate(comm, ncclFuncAllReduce, algorithm, nBytes); +} +int RingProtocol(ncclComm* comm, ncclFunc_t func, size_t perRank, ncclDataType_t dt = ncclFloat32) { + return ProtocolAfterUpdate(comm, func, NCCL_ALGO_RING, perRank, dt); +} +} // namespace + +// AICOMRCCL-1756: on 8+ rank single-node gfx1250 comms Tree AllReduce stays on LL up to +// nRanks x 32 KiB and leaves it one byte later. 4 ranks, two nodes, or no arch table +// keep the model's 128 KiB LL range. Holds with or without ENABLE_LL128. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250TreeAllReduceLLScalesWithRanks) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250TreeAllReduceLLScalesWithRanks", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + for (int nRanks : {8, 64}) { + ncclComm* comm = MakeGfx1250ProtocolComm(nRanks); + const size_t cap = size_t(nRanks) * 32 * kKiB; + EXPECT_EQ(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, cap)) << nRanks; + EXPECT_NE(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, cap + 1)) << nRanks; + EXPECT_NE(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_RING, cap)) << "Ring keeps the model range, " << nRanks; + DeleteCommWithArch(comm); + } + ncclComm* comm = MakeGfx1250ProtocolComm(4); + EXPECT_EQ(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, 131072)); + EXPECT_NE(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, 131073)); + comm->nRanks = 8; + comm->nNodes = 2; + EXPECT_NE(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, 131073)) << "two nodes"; + comm->nNodes = 1; + comm->archThresholds = nullptr; + EXPECT_NE(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, 131073)) << "no arch table"; + DeleteCommWithArch(comm); + }); +} + +// The LL128 overrides sit inside rcclUpdateCollectiveProtocol's ENABLE_LL128 arm, so +// these only compile where that arm does (see the IsArchSupportedForFunc_LL128_* pair +// above for the same split). +#if defined(ENABLE_LL128) +// AICOMRCCL-2874: on gfx1250 Tree LL128 beats Tree Simple at every size, so a Tree +// AllReduce far past the model's LL128 ceiling stays on LL128. Ring at the same size, +// and Tree on a two-node comm, are the controls: they still drop to Simple. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250TreeAllReduceKeepsLL128) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250TreeAllReduceKeepsLL128", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + ncclComm* comm = MakeGfx1250ProtocolComm(16); + const size_t big = size_t(16) << 30; + EXPECT_EQ(NCCL_PROTO_LL128, ArProtocol(comm, NCCL_ALGO_TREE, big)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, NCCL_ALGO_RING, big)); + EXPECT_EQ(NCCL_PROTO_LL, ArProtocol(comm, NCCL_ALGO_TREE, 1024)) << "LL range still wins small sizes"; + comm->nNodes = 2; + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, NCCL_ALGO_TREE, big)) << "two nodes"; + DeleteCommWithArch(comm); + }); +} + +// AICOMRCCL-2875: Ring AllReduce keeps LL128 up to nRanks x 32 MiB (inclusive), past the +// model's 16 MiB ceiling, and returns to Simple one byte later. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250RingAllReduceLL128ScalesWithRanks) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250RingAllReduceLL128ScalesWithRanks", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + // Not a rank-scaled entry: it applies below rankScaledMinRanks too. + for (int nRanks : {4, 8, 16, 64}) { + ncclComm* comm = MakeGfx1250ProtocolComm(nRanks); + const size_t cap = size_t(nRanks) * 32 * kMiB; + EXPECT_EQ(NCCL_PROTO_LL128, ArProtocol(comm, NCCL_ALGO_RING, cap)) << nRanks; + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, NCCL_ALGO_RING, cap + 1)) << nRanks; + DeleteCommWithArch(comm); + } + // Two ranks with a 128 MiB model ceiling: nRanks x 32 MiB (64 MiB) sits below it, + // so the model ceiling is the one that applies. + constexpr uint64_t kModelCeiling = 128ULL << 20; + ncclComm* comm = MakeGfx1250ProtocolComm(2); + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MAX_IDX] = kModelCeiling; + EXPECT_EQ(NCCL_PROTO_LL128, ArProtocol(comm, NCCL_ALGO_RING, kModelCeiling)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, NCCL_ALGO_RING, kModelCeiling + 1)); + DeleteCommWithArch(comm); + }); +} + +// Other arches keep the model ceiling for both algorithms. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_NonGfx1250AllReduceKeepsTableLL128Ceiling) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_NonGfx1250AllReduceKeepsTableLL128Ceiling", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + // gfx1250's ranges and rank count on another arch, which reaches the same arm via nNodes >= 2. + ncclComm* comm = MakeCommWithArch("gfx90a"); + comm->nNodes = 2; + comm->nRanks = 16; + comm->topo->ll128Enabled = true; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL][RCCL_PROTOCOL_MAX_IDX] = 131072; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MIN_IDX] = 131072; + comm->minMaxLLRange[RCCL_AR_TUNABLE][NCCL_PROTO_LL128][RCCL_PROTOCOL_MAX_IDX] = kGfx1250ArLL128Max; + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, NCCL_ALGO_TREE, kGfx1250ArLL128Max + 1)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, NCCL_ALGO_RING, kGfx1250ArLL128Max + 1)); + EXPECT_EQ(NCCL_PROTO_LL128, ArProtocol(comm, NCCL_ALGO_RING, kGfx1250ArLL128Max)); + DeleteCommWithArch(comm); + }); +} + +// Tree AllReduce takes LL128 right above the scaled LL range. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250TreeAllReduceLL128AboveLLRange) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250TreeAllReduceLL128AboveLLRange", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + ncclComm* comm = MakeGfx1250ProtocolComm(8); + EXPECT_EQ(NCCL_PROTO_LL128, ArProtocol(comm, NCCL_ALGO_TREE, 8 * 32 * kKiB + 1)); + DeleteCommWithArch(comm); + }); +} + +// AllGather Ring keeps LL128 to 64 MiB per rank on 8+ rank single-node gfx1250 comms +// (model: 16 MiB); two-node and 4-rank comms keep the model. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250AllGatherRingLL128To64MiBPerRank) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250AllGatherRingLL128To64MiBPerRank", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + ncclComm* comm = MakeGfx1250ProtocolComm(16); + EXPECT_EQ(NCCL_PROTO_LL128, RingProtocol(comm, ncclFuncAllGather, 64 * kMiB)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, RingProtocol(comm, ncclFuncAllGather, 64 * kMiB + 16)); + comm->nNodes = 2; + EXPECT_EQ(NCCL_PROTO_SIMPLE, RingProtocol(comm, ncclFuncAllGather, 16 * kMiB + 16)) << "two nodes"; + comm->nNodes = 1; + comm->nRanks = 4; + EXPECT_EQ(NCCL_PROTO_LL128, RingProtocol(comm, ncclFuncAllGather, 16 * kMiB)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, RingProtocol(comm, ncclFuncAllGather, 16 * kMiB + 16)); + DeleteCommWithArch(comm); + }); +} + +// ReduceScatter Ring on 8+ rank single-node gfx1250 comms: LL to 512 KiB per rank, LL128 +// to 8 MiB per rank (32 MiB for 1-byte types), Simple above. Two-node and 4-rank comms +// keep the model ranges. +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250ReduceScatterRingRanges) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250ReduceScatterRingRanges", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + ncclComm* comm = MakeGfx1250ProtocolComm(8); + EXPECT_EQ(NCCL_PROTO_LL, RingProtocol(comm, ncclFuncReduceScatter, 512 * kKiB)); + EXPECT_EQ(NCCL_PROTO_LL128, RingProtocol(comm, ncclFuncReduceScatter, 512 * kKiB + 16)); + EXPECT_EQ(NCCL_PROTO_LL128, RingProtocol(comm, ncclFuncReduceScatter, 8 * kMiB)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, RingProtocol(comm, ncclFuncReduceScatter, 8 * kMiB + 16)); + EXPECT_EQ(NCCL_PROTO_LL128, RingProtocol(comm, ncclFuncReduceScatter, 32 * kMiB, ncclFloat8e4m3)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, RingProtocol(comm, ncclFuncReduceScatter, 32 * kMiB + 16, ncclFloat8e4m3)); + comm->nNodes = 2; + EXPECT_EQ(NCCL_PROTO_LL, RingProtocol(comm, ncclFuncReduceScatter, 1 * kMiB)) << "two nodes"; + comm->nNodes = 1; + comm->nRanks = 4; + EXPECT_EQ(NCCL_PROTO_LL, RingProtocol(comm, ncclFuncReduceScatter, 1 * kMiB)); + EXPECT_EQ(NCCL_PROTO_LL128, RingProtocol(comm, ncclFuncReduceScatter, 4 * kMiB)); + EXPECT_EQ(NCCL_PROTO_SIMPLE, RingProtocol(comm, ncclFuncReduceScatter, 4 * kMiB + 16)); + DeleteCommWithArch(comm); + }); +} +#else +TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_Gfx1250AllReduceLL128_DisabledAtCompileTime) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveProtocol_Gfx1250AllReduceLL128_DisabledAtCompileTime", + []() { + SetMicroEnvAbsent("NCCL_PROTO"); + ncclComm* comm = MakeGfx1250ProtocolComm(16); + for (int algorithm : {NCCL_ALGO_TREE, NCCL_ALGO_RING}) { + EXPECT_EQ(NCCL_PROTO_SIMPLE, ArProtocol(comm, algorithm, 64 * kMiB)) << "algorithm=" << algorithm; + } + DeleteCommWithArch(comm); + }); +} +#endif // ENABLE_LL128 + TEST(WrapMicrotestIsolated, UpdateCollectiveProtocol_UndefinedTuningWarnsOnceForSupportedArch) { RUN_ISOLATED_TEST( "Wrap_UpdateCollectiveProtocol_UndefinedTuningWarnsOnceForSupportedArch", @@ -4733,6 +4997,648 @@ TEST(WrapMicrotestIsolated, SelectAllGather_CeForceScratchNotChosenWhenProbeSays }); } +namespace { +// rccl-UnitTestsMicro links no real arch tables (the host fake rcclGetArchThresholds +// returns null), so each case hands the comm the table it wants to exercise. +// Every other CE-Scratch gate is armed: scratch probe true, unregistered windows, +// no sysmem, receive fits ddaScratch, and ceNonRegMax[AG] = 0 so the window +// cannot fire on its own. +rcclArchThresholds MakeCeScratchTable(bool forceCeScratch) { + rcclArchThresholds table{}; + table.forceCeScratch = forceCeScratch; + return table; +} + +ncclComm* MakeCeScratchComm(const char* arch, const rcclArchThresholds* table, uint8_t* scratch, size_t bytes) { + ncclComm* comm = MakeCommWithArch(arch); + comm->nRanks = 1; + comm->nNodes = 1; + comm->archThresholds = table; + comm->ddaScratch = scratch; + comm->ddaScratchBytes = bytes; + return comm; +} + +int SelectAllGatherAlgo(ncclComm* comm) { + rcclCollDecision decision{}; + EXPECT_EQ(ncclSuccess, rcclSelectAllGather(comm, nullptr, nullptr, /*sendcount=*/8, ncclFloat32, + /*stream=*/nullptr, /*query=*/true, /*graphCapturingHint=*/false, + &decision)); + return decision.algo; +} + +int SelectAlltoAllAlgo(ncclComm* comm) { + rcclCollDecision decision{}; + EXPECT_EQ(ncclSuccess, rcclSelectAlltoAll(comm, nullptr, nullptr, /*count=*/8, ncclFloat32, /*stream=*/nullptr, + /*query=*/true, /*graphCapturingHint=*/false, &decision)); + return decision.algo; +} +} // namespace + +// AICOMRCCL-2876: an arch table with forceCeScratch = false (gfx1250) leaves CE-Scratch +// unforced under the default RCCL_FORCE_CE=1. Flipping only the table entry is the +// control: it proves every other gate is armed, so the entry alone keeps CE-Scratch off. +TEST(WrapMicrotestIsolated, SelectAllGather_TableWithoutForceCeScratchDoesNotForce) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllGather_TableWithoutForceCeScratchDoesNotForce", + []() { + ScopedHook ceScratch(g_ceScratchAvailable, [](struct ncclComm*, ncclFunc_t, int, ncclDataType_t, + ncclSymRegType_t) { return true; }); + rcclArchThresholds table = MakeCeScratchTable(/*forceCeScratch=*/false); + uint8_t scratch[64]; + ncclComm* comm = MakeCeScratchComm("gfx1250", &table, scratch, sizeof(scratch)); + EXPECT_EQ(NCCL_ALGO_RING, SelectAllGatherAlgo(comm)); + table.forceCeScratch = true; + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAllGatherAlgo(comm)); + DeleteCommWithArch(comm); + }); +} + +// CE-Scratch is the one zero-CU path for unregistered buffers, so a comm-level +// NCCL_CTA_POLICY_ZERO keeps it forced even when the table does not. An explicit +// FORCE_CE=0 still wins. +TEST(WrapMicrotestIsolated, SelectAllGather_CtaZeroKeepsCeScratch) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllGather_CtaZeroKeepsCeScratch", + []() { + ScopedHook ceScratch(g_ceScratchAvailable, [](struct ncclComm*, ncclFunc_t, int, ncclDataType_t, + ncclSymRegType_t) { return true; }); + const rcclArchThresholds table = MakeCeScratchTable(/*forceCeScratch=*/false); + uint8_t scratch[64]; + ncclComm* comm = MakeCeScratchComm("gfx1250", &table, scratch, sizeof(scratch)); + comm->config.CTAPolicy = NCCL_CTA_POLICY_ZERO; + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAllGatherAlgo(comm)); + { + ScopedHook forceCe(g_paramForceCe, []() { return int64_t(0); }); + EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAllGatherAlgo(comm)); + } + DeleteCommWithArch(comm); + }); +} + +// A table with forceCeScratch = true (gfx942, gfx950) forces CE-Scratch, as does a +// comm with no table (unknown arch or RCCL_IGNORE_ARCH_TABLE): MI300/MI355 keep the +// old behaviour. An explicit FORCE_CE=0 turns it off in both cases. +TEST(WrapMicrotestIsolated, SelectAllGather_ForceCeScratchTableOrNoTableForces) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllGather_ForceCeScratchTableOrNoTableForces", + []() { + ScopedHook ceScratch(g_ceScratchAvailable, [](struct ncclComm*, ncclFunc_t, int, ncclDataType_t, + ncclSymRegType_t) { return true; }); + const rcclArchThresholds table = MakeCeScratchTable(/*forceCeScratch=*/true); + for (const rcclArchThresholds* t : {&table, static_cast(nullptr)}) { + for (const char* arch : {"gfx942", "gfx950"}) { + uint8_t scratch[64]; + ncclComm* comm = MakeCeScratchComm(arch, t, scratch, sizeof(scratch)); + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAllGatherAlgo(comm)) + << arch << (t ? " table" : " no table"); + { + ScopedHook forceCe(g_paramForceCe, []() { return int64_t(0); }); + EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAllGatherAlgo(comm)) + << arch << (t ? " table" : " no table"); + } + DeleteCommWithArch(comm); + } + } + }); +} + +// The AlltoAll selector reads the same resolver. +TEST(WrapMicrotestIsolated, SelectAlltoAll_CeScratchFollowsForceCeScratchResolver) { + RUN_ISOLATED_TEST( + "Wrap_SelectAlltoAll_CeScratchFollowsForceCeScratchResolver", + []() { + ScopedHook ceScratch(g_ceScratchAvailable, [](struct ncclComm*, ncclFunc_t, int, ncclDataType_t, + ncclSymRegType_t) { return true; }); + rcclArchThresholds table = MakeCeScratchTable(/*forceCeScratch=*/false); + uint8_t scratch[64]; + ncclComm* comm = MakeCeScratchComm("gfx1250", &table, scratch, sizeof(scratch)); + EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAlltoAllAlgo(comm)); + comm->config.CTAPolicy = NCCL_CTA_POLICY_ZERO; + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAlltoAllAlgo(comm)); + comm->config.CTAPolicy = 0; + table.forceCeScratch = true; + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAlltoAllAlgo(comm)); + { + ScopedHook forceCe(g_paramForceCe, []() { return int64_t(0); }); + EXPECT_NE((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAlltoAllAlgo(comm)); + } + DeleteCommWithArch(comm); + }); +} + +// The LL128 two-shot AllReduce ceiling (AICOMRCCL-2875) follows the AllReduce LL128 +// tier cap where the arch table runs that tier; a table without it, or no table, +// keeps RCCL_DDA_LL128_TWOSHOT_THRESHOLD's 64 MiB default. +TEST(WrapMicrotestIsolated, DdaLL128TwoShotThreshold_FollowsLL128TierElseEnvDefault) { + RUN_ISOLATED_TEST( + "Wrap_DdaLL128TwoShotThreshold_FollowsLL128TierElseEnvDefault", + []() { + SetMicroEnvAbsent("RCCL_DDA_LL128_TWOSHOT_THRESHOLD"); + rcclArchThresholds table = MakeRankScaledTable(); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = &table; + comm->nRanks = 8; + EXPECT_EQ(256 * kMiB, rcclDdaLL128TwoShotThreshold(comm)); + table.ddaLL128Max[ncclFuncAllReduce] = 0; + EXPECT_EQ(64 * kMiB, rcclDdaLL128TwoShotThreshold(comm)); + comm->archThresholds = nullptr; + EXPECT_EQ(64 * kMiB, rcclDdaLL128TwoShotThreshold(comm)) << "no table"; + DeleteCommWithArch(comm); + }); +} + +// A user-set RCCL_DDA_LL128_TWOSHOT_THRESHOLD overrides the tier cap. +TEST(WrapMicrotestIsolated, DdaLL128TwoShotThreshold_EnvOverridesTable) { + RUN_ISOLATED_TEST( + "Wrap_DdaLL128TwoShotThreshold_EnvOverridesTable", + []() { + SetMicroEnv("RCCL_DDA_LL128_TWOSHOT_THRESHOLD", "33554432"); + g_loadParam = ForceParam("RCCL_DDA_LL128_TWOSHOT_THRESHOLD", int64_t(32) << 20); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = Gfx1250Table(); + comm->nRanks = 8; + EXPECT_EQ(32 * kMiB, rcclDdaLL128TwoShotThreshold(comm)); + DeleteCommWithArch(comm); + }); +} + +// An empty value is not a setting: the param loader reads it as unset, so the +// resolver must too, and the tier cap stays in force. +TEST(WrapMicrotestIsolated, DdaLL128TwoShotThreshold_EmptyEnvIsUnset) { + RUN_ISOLATED_TEST( + "Wrap_DdaLL128TwoShotThreshold_EmptyEnvIsUnset", + []() { + SetMicroEnv("RCCL_DDA_LL128_TWOSHOT_THRESHOLD", ""); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = Gfx1250Table(); + comm->nRanks = 8; + EXPECT_EQ(256 * kMiB, rcclDdaLL128TwoShotThreshold(comm)); + DeleteCommWithArch(comm); + }); +} + +// =========================================================================== +// AICOMRCCL-1756: rank-scaled gfx1250 DDA tiers, Tree for small AllReduce, +// Ring protocol ranges and the ReduceScatter channel cap. +// =========================================================================== + +// Below rankScaledMinRanks the totals apply; from it, nRanks x the per-rank entry, +// clipped to rankScaledMaxBytes. Tiers without a per-rank entry keep their totals. +TEST(WrapMicrotestIsolated, DdaTierCaps_RankScaledFromMinRanks) { + RUN_ISOLATED_TEST( + "Wrap_DdaTierCaps_RankScaledFromMinRanks", + []() { + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = Gfx1250Table(); + comm->nRanks = 7; + EXPECT_EQ(32 * kMiB, rcclDdaLLThreshold(comm, ncclFuncAllReduce)); + EXPECT_EQ(128 * kKiB, rcclDdaLLThreshold(comm, ncclFuncAllGather)); + EXPECT_EQ(2 * kMiB, rcclDdaLL128Threshold(comm, ncclFuncReduceScatter)); + EXPECT_EQ(0u, rcclDdaVmmThreshold(comm, ncclFuncReduceScatter)); + comm->nRanks = 8; + EXPECT_EQ(8 * kMiB, rcclDdaLLThreshold(comm, ncclFuncAllReduce)); + EXPECT_EQ(128 * kKiB, rcclDdaLLThreshold(comm, ncclFuncAllGather)); + EXPECT_EQ(2 * kMiB, rcclDdaLLThreshold(comm, ncclFuncReduceScatter)); + EXPECT_EQ(128 * kMiB, rcclDdaLL128Threshold(comm, ncclFuncAllGather)); + EXPECT_EQ(4 * kMiB, rcclDdaLL128Threshold(comm, ncclFuncReduceScatter)); + EXPECT_EQ(16 * kMiB, rcclDdaVmmThreshold(comm, ncclFuncReduceScatter)); + EXPECT_EQ(256 * kMiB, rcclDdaLL128Threshold(comm, ncclFuncAllReduce)) << "no per-rank entry"; + EXPECT_EQ(16 * kMiB, rcclDdaVmmThreshold(comm, ncclFuncAllReduce)) << "no per-rank entry"; + comm->nRanks = 32; + EXPECT_EQ(256 * kMiB, rcclDdaLL128Threshold(comm, ncclFuncAllGather)) << "clipped"; + EXPECT_EQ(64 * kMiB, rcclDdaVmmThreshold(comm, ncclFuncReduceScatter)); + // The entry gate admits the widest tier, so it scales too. + EXPECT_EQ(64 * kMiB, rcclDdaEntryThreshold(comm, ncclFuncReduceScatter)); + DeleteCommWithArch(comm); + }); +} + +// RCCL_DDA_LL_THRESHOLD still wins over the rank-scaled cap. +TEST(WrapMicrotestIsolated, DdaTierCaps_EnvWinsOverRankScaling) { + RUN_ISOLATED_TEST( + "Wrap_DdaTierCaps_EnvWinsOverRankScaling", + []() { + g_loadParam = ForceParam("RCCL_DDA_LL_THRESHOLD", int64_t(65536)); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = Gfx1250Table(); + comm->nRanks = 16; + EXPECT_EQ(65536u, rcclDdaLLThreshold(comm, ncclFuncAllReduce)); + EXPECT_EQ(65536u, rcclDdaLLThreshold(comm, ncclFuncReduceScatter)); + DeleteCommWithArch(comm); + }); +} + +// On rank-scaled comms the AllReduce LL two-shot runs to the LL tier cap; smaller +// comms, and tables that do not scale the AllReduce LL tier, keep the 2 MiB default. +TEST(WrapMicrotestIsolated, DdaLLTwoShotThreshold_FollowsRankScaledLLTier) { + RUN_ISOLATED_TEST( + "Wrap_DdaLLTwoShotThreshold_FollowsRankScaledLLTier", + []() { + SetMicroEnvAbsent("RCCL_DDA_LL_TWOSHOT_THRESHOLD"); + rcclArchThresholds table = MakeRankScaledTable(); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = &table; + comm->nRanks = 16; + EXPECT_EQ(16 * kMiB, rcclDdaLLTwoShotThreshold(comm)); + comm->nRanks = 4; + EXPECT_EQ(2 * kMiB, rcclDdaLLTwoShotThreshold(comm)); + comm->nRanks = 16; + table.ddaLLMaxPerRank[ncclFuncAllReduce] = 0; + EXPECT_EQ(2 * kMiB, rcclDdaLLTwoShotThreshold(comm)); + comm->archThresholds = nullptr; + EXPECT_EQ(2 * kMiB, rcclDdaLLTwoShotThreshold(comm)) << "no table"; + DeleteCommWithArch(comm); + }); +} + +TEST(WrapMicrotestIsolated, DdaLLTwoShotThreshold_EnvOverridesRankScaling) { + RUN_ISOLATED_TEST( + "Wrap_DdaLLTwoShotThreshold_EnvOverridesRankScaling", + []() { + SetMicroEnv("RCCL_DDA_LL_TWOSHOT_THRESHOLD", "4194304"); + g_loadParam = ForceParam("RCCL_DDA_LL_TWOSHOT_THRESHOLD", int64_t(4) << 20); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = Gfx1250Table(); + comm->nRanks = 16; + EXPECT_EQ(4 * kMiB, rcclDdaLLTwoShotThreshold(comm)); + DeleteCommWithArch(comm); + }); +} + +// VMM blocks: clamp(2048 / nRanks, 32, 128), never above the init cap; small comms, +// no budget or no table keep the init cap. +TEST(WrapMicrotestIsolated, DdaFabricVmmBlocks_BudgetPerRank) { + RUN_ISOLATED_TEST( + "Wrap_DdaFabricVmmBlocks_BudgetPerRank", + []() { + rcclArchThresholds table = MakeRankScaledTable(); + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->archThresholds = &table; + comm->ddaFabricMaxBlocks = 192; + const std::pair cases[] = {{4, 192}, {8, 128}, {16, 128}, {32, 64}, {64, 32}, {128, 32}}; + for (const auto& c : cases) { + comm->nRanks = c.first; + EXPECT_EQ(c.second, rcclDdaFabricVmmBlocks(comm)) << c.first; + } + comm->nRanks = 8; + comm->ddaFabricMaxBlocks = 48; + EXPECT_EQ(48, rcclDdaFabricVmmBlocks(comm)) << "init cap still bounds it"; + comm->ddaFabricMaxBlocks = 192; + table.ddaVmmBlockBudget = 0; + EXPECT_EQ(192, rcclDdaFabricVmmBlocks(comm)); + comm->archThresholds = nullptr; + EXPECT_EQ(192, rcclDdaFabricVmmBlocks(comm)); + DeleteCommWithArch(comm); + }); +} + +namespace { +int SelectAllReduceDdaAlgo(ncclComm* comm, size_t bytes) { + rcclCollDecision decision{}; + EXPECT_EQ(ncclSuccess, rcclSelectAllReduce(comm, nullptr, nullptr, bytes / sizeof(float), ncclFloat32, ncclSum, + /*stream=*/nullptr, /*query=*/true, /*graphCapturingHint=*/false, + &decision)); + return decision.algo; +} + +template +struct HookSig; +template +struct HookSig> { + using type = Sig; +}; +template +using HookFor = ScopedHook::type>; + +// AllReduce DDA path taken and its LL tier eligible; ArVmmEligible adds the VMM tier. +struct ArDdaHooks { + HookFor shouldTakeDda{ + g_allReduceShouldTakeDdaPath, [](const struct ncclComm*, size_t, ncclDataType_t, bool, bool, bool) { return true; }}; + HookFor llEligible{ + g_allReduceDdaFabricLLEligible, [](ncclComm*, const void*, void*, size_t, ncclDataType_t, ncclRedOp_t) { return true; }}; +}; +struct ArVmmEligible { + HookFor vmmEligible{ + g_allReduceDdaFabricEligible, [](ncclComm*, const void*, void*, size_t, ncclDataType_t, ncclRedOp_t) { return true; }}; +}; +ncclComm* MakeArDdaComm(int nRanks) { + ncclComm* comm = MakeCommWithArch("gfx1250"); + comm->nRanks = nRanks; + comm->archThresholds = Gfx1250Table(); + return comm; +} +} // namespace + +// From 32 ranks, DDA VMM claims (2 MiB, nRanks x 512 KiB] ahead of LL; at and below +// 2 MiB, past the window, or on a 16-rank comm, LL keeps it. +TEST(WrapMicrotestIsolated, SelectAllReduce_Gfx1250VmmFirstWindow) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllReduce_Gfx1250VmmFirstWindow", + []() { + ArDdaHooks hooks; + ArVmmEligible vmmHook; + ncclComm* comm = MakeArDdaComm(32); + const int vmm = (int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_VMM; + const int ll = (int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_LL; + EXPECT_EQ(ll, SelectAllReduceDdaAlgo(comm, 2 * kMiB)); + EXPECT_EQ(vmm, SelectAllReduceDdaAlgo(comm, 2 * kMiB + 16)); + EXPECT_EQ(vmm, SelectAllReduceDdaAlgo(comm, 16 * kMiB)); + EXPECT_EQ(ll, SelectAllReduceDdaAlgo(comm, 16 * kMiB + 16)); + comm->nRanks = 64; + EXPECT_EQ(vmm, SelectAllReduceDdaAlgo(comm, 24 * kMiB)) << "past ddaVmmMax (16 MiB) inside the window"; + comm->nRanks = 16; + EXPECT_EQ(ll, SelectAllReduceDdaAlgo(comm, 4 * kMiB)); + DeleteCommWithArch(comm); + }); +} + +// Registered receive windows keep the normal tier order (the window was measured on +// unregistered buffers). +TEST(WrapMicrotestIsolated, SelectAllReduce_Gfx1250VmmFirstUnregisteredOnly) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllReduce_Gfx1250VmmFirstUnregisteredOnly", + []() { + ArDdaHooks hooks; + ArVmmEligible vmmHook; + ScopedHook regType(g_getSymRegType, [](struct ncclDevrWindow*, struct ncclDevrWindow*, ncclSymRegType_t* out) { + *out = ncclSymSendNonregRecvReg; + return ncclSuccess; + }); + ncclComm* comm = MakeArDdaComm(32); + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_LL, SelectAllReduceDdaAlgo(comm, 8 * kMiB)); + DeleteCommWithArch(comm); + }); +} + +// The window steps aside when the user sets RCCL_DDA_LL_THRESHOLD ... +TEST(WrapMicrotestIsolated, SelectAllReduce_Gfx1250VmmFirstYieldsToLLThresholdEnv) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllReduce_Gfx1250VmmFirstYieldsToLLThresholdEnv", + []() { + g_loadParam = ForceParam("RCCL_DDA_LL_THRESHOLD", int64_t(64) << 20); + ArDdaHooks hooks; + ArVmmEligible vmmHook; + ncclComm* comm = MakeArDdaComm(64); + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_LL, SelectAllReduceDdaAlgo(comm, 8 * kMiB)); + DeleteCommWithArch(comm); + }); +} + +// ... or RCCL_DDA_THRESHOLD. +TEST(WrapMicrotestIsolated, SelectAllReduce_Gfx1250VmmFirstYieldsToVmmThresholdEnv) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllReduce_Gfx1250VmmFirstYieldsToVmmThresholdEnv", + []() { + g_loadParam = ForceParam("RCCL_DDA_THRESHOLD", int64_t(64) << 20); + ArDdaHooks hooks; + ArVmmEligible vmmHook; + ncclComm* comm = MakeArDdaComm(64); + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_LL, SelectAllReduceDdaAlgo(comm, 8 * kMiB)); + DeleteCommWithArch(comm); + }); +} + +TEST(WrapMicrotestIsolated, SelectAllReduce_Gfx1250VmmFirstNeedsVmmEligible) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllReduce_Gfx1250VmmFirstNeedsVmmEligible", + []() { + ArDdaHooks hooks; // g_allReduceDdaFabricEligible left at its default (false) + ncclComm* comm = MakeArDdaComm(64); + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_LL, SelectAllReduceDdaAlgo(comm, 8 * kMiB)); + DeleteCommWithArch(comm); + }); +} + +// The ReduceScatter selector passes nRanks to the VMM cap: VMM, newly on for gfx1250, +// runs to nRanks x 2 MiB on a 16-rank comm and not past it; 4 ranks keep it off. +TEST(WrapMicrotestIsolated, SelectReduceScatter_Gfx1250RankScaledVmm) { + RUN_ISOLATED_TEST( + "Wrap_SelectReduceScatter_Gfx1250RankScaledVmm", + []() { + ScopedHook vmmEligible(g_reduceScatterDdaFabricEligible, + [](ncclComm*, const void*, void*, size_t, ncclDataType_t, ncclRedOp_t) { return true; }); + ncclComm* comm = MakeArDdaComm(16); + auto algo = [&](size_t totalBytes) { + rcclCollDecision decision{}; + EXPECT_EQ(ncclSuccess, rcclSelectReduceScatter(comm, nullptr, nullptr, + totalBytes / sizeof(float) / (size_t)comm->nRanks, + ncclFloat32, ncclSum, /*query=*/true, &decision)); + return decision.algo; + }; + const int vmm = (int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_VMM; + EXPECT_EQ(vmm, algo(32 * kMiB)); + EXPECT_NE(vmm, algo(32 * kMiB + 16 * 16 * 4)); + comm->nRanks = 4; + EXPECT_NE(vmm, algo(1 * kMiB)); + DeleteCommWithArch(comm); + }); +} + +// The AllGather selector passes nRanks to the LL128 cap: 256 MiB on a 16-rank comm, +// 128 MiB (the total) on a 4-rank one. +TEST(WrapMicrotestIsolated, SelectAllGather_Gfx1250RankScaledLL128) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllGather_Gfx1250RankScaledLL128", + []() { + ScopedHook ll128Eligible(g_allGatherDdaFabricLL128Eligible, + [](ncclComm*, const void*, void*, size_t, ncclDataType_t) { return true; }); + ncclComm* comm = MakeArDdaComm(16); + auto algo = [&](size_t totalBytes) { + rcclCollDecision decision{}; + EXPECT_EQ(ncclSuccess, rcclSelectAllGather(comm, nullptr, nullptr, + totalBytes / sizeof(float) / (size_t)comm->nRanks, ncclFloat32, + /*stream=*/nullptr, /*query=*/true, + /*graphCapturingHint=*/false, &decision)); + return decision.algo; + }; + const int ll128 = (int)rcclAddonAlgos_t::RCCL_DDA_FABRIC_LL128; + EXPECT_EQ(ll128, algo(256 * kMiB)); + EXPECT_NE(ll128, algo(256 * kMiB + 16 * 16)); + comm->nRanks = 4; + EXPECT_EQ(ll128, algo(128 * kMiB)); + EXPECT_NE(ll128, algo(128 * kMiB + 4 * 16)); + DeleteCommWithArch(comm); + }); +} + +// An explicit RCCL_FORCE_CE=1 still forces CE-Scratch where the table does not. +TEST(WrapMicrotestIsolated, SelectAllGather_ExplicitForceCeForcesCeScratch) { + RUN_ISOLATED_TEST( + "Wrap_SelectAllGather_ExplicitForceCeForcesCeScratch", + []() { + SetMicroEnv("RCCL_FORCE_CE", "1"); + ScopedHook ceScratch(g_ceScratchAvailable, [](struct ncclComm*, ncclFunc_t, int, ncclDataType_t, + ncclSymRegType_t) { return true; }); + const rcclArchThresholds table = MakeCeScratchTable(/*forceCeScratch=*/false); + uint8_t scratch[64]; + ncclComm* comm = MakeCeScratchComm("gfx1250", &table, scratch, sizeof(scratch)); + EXPECT_EQ((int)rcclAddonAlgos_t::RCCL_CE_SCRATCH, SelectAllGatherAlgo(comm)); + DeleteCommWithArch(comm); + }); +} + +namespace { +// Cost table offering every algorithm/protocol (Ring cheapest), as topoGetAlgoInfo sees it. +struct CostTable { + float t[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS]; + CostTable() { + for (auto& row : t) + for (float& v : row) v = 2.0f; + } +}; + +int AlgoAfterUpdate(ncclComm* comm, ncclFunc_t func, size_t nBytes, CostTable& table) { + ncclTaskColl info{}; + info.func = func; + info.algorithm = NCCL_ALGO_RING; + info.protocol = NCCL_PROTO_LL128; + rcclUpdateCollectiveAlgorithm(comm, nBytes, &info, table.t); + return info.algorithm; +} + +// Single-node comm on `arch`; gfx1250 gets the arch table. +ncclComm* MakeTuningComm(int nRanks, const char* arch = "gfx1250") { + ncclComm* comm = MakeCommWithArch(arch); + comm->nRanks = nRanks; + comm->nNodes = 1; + if (std::strcmp(arch, "gfx1250") == 0) comm->archThresholds = Gfx1250Table(); + return comm; +} +} // namespace + +// AICOMRCCL-1756: gfx1250 AllReduce on 8+ ranks runs Tree up to nRanks x 1 MiB. +TEST(WrapMicrotestIsolated, UpdateCollectiveAlgorithm_Gfx1250SmallAllReduceTakesTree) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveAlgorithm_Gfx1250SmallAllReduceTakesTree", + []() { + SetMicroEnvAbsent("NCCL_ALGO"); + SetMicroEnvAbsent("NCCL_PROTO"); + CostTable table; + for (int nRanks : {8, 32}) { + ncclComm* comm = MakeTuningComm(nRanks); + const size_t cap = (size_t)nRanks * kMiB; + EXPECT_EQ(NCCL_ALGO_TREE, AlgoAfterUpdate(comm, ncclFuncAllReduce, cap, table)) << nRanks; + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, cap + 1, table)) << nRanks; + DeleteCommWithArch(comm); + } + }); +} + +// Each guard on its own: below 8 ranks, more than one node, no arch table, another +// arch, another collective, or Tree LL / LL128 missing from the cost table all keep Ring. +TEST(WrapMicrotestIsolated, UpdateCollectiveAlgorithm_TreeOnlyWhenEveryGuardHolds) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveAlgorithm_TreeOnlyWhenEveryGuardHolds", + []() { + SetMicroEnvAbsent("NCCL_ALGO"); + SetMicroEnvAbsent("NCCL_PROTO"); + CostTable table; + ncclComm* comm = MakeTuningComm(4); + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)) << "4 ranks"; + comm->nRanks = 8; + EXPECT_EQ(NCCL_ALGO_TREE, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)) << "control"; + comm->nNodes = 2; + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)) << "2 nodes"; + comm->nNodes = 1; + comm->archThresholds = nullptr; + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)) << "no arch table"; + comm->archThresholds = Gfx1250Table(); + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllGather, 1024, table)) << "AllGather"; + table.t[NCCL_ALGO_TREE][NCCL_PROTO_LL] = NCCL_ALGO_PROTO_IGNORE; + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)) << "no Tree LL"; + table.t[NCCL_ALGO_TREE][NCCL_PROTO_LL] = 2.0f; + table.t[NCCL_ALGO_TREE][NCCL_PROTO_LL128] = NCCL_ALGO_PROTO_IGNORE; + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)) << "no Tree LL128"; + DeleteCommWithArch(comm); + CostTable full; + comm = MakeTuningComm(8, "gfx950"); + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, full)) << "gfx950"; + DeleteCommWithArch(comm); + }); +} + +TEST(WrapMicrotestIsolated, UpdateCollectiveAlgorithm_NcclAlgoEnvKeepsRing) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveAlgorithm_NcclAlgoEnvKeepsRing", + []() { + SetMicroEnv("NCCL_ALGO", "Ring"); + CostTable table; + ncclComm* comm = MakeTuningComm(8); + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)); + DeleteCommWithArch(comm); + }); +} + +// With NCCL_PROTO set the protocol hook leaves Ring's protocol in place, so the +// switch to Tree must not happen either. +TEST(WrapMicrotestIsolated, UpdateCollectiveAlgorithm_NcclProtoEnvKeepsRing) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveAlgorithm_NcclProtoEnvKeepsRing", + []() { + SetMicroEnvAbsent("NCCL_ALGO"); + SetMicroEnv("NCCL_PROTO", "Simple,LL,LL128"); + CostTable table; + ncclComm* comm = MakeTuningComm(8); + EXPECT_EQ(NCCL_ALGO_RING, AlgoAfterUpdate(comm, ncclFuncAllReduce, 1024, table)); + DeleteCommWithArch(comm); + }); +} + +namespace { +int CappedChannels(ncclComm* comm, ncclFunc_t func, int proto, size_t perRank) { + ncclTaskColl info{}; + info.func = func; + info.algorithm = NCCL_ALGO_RING; + info.protocol = proto; + int nc = 192; + rcclUpdateCollectiveChannels(comm, perRank * (size_t)comm->nRanks, &info, nc); + return nc; +} +} // namespace + +// Ring Simple ReduceScatter between 8 and 32 MiB per rank runs on 128 channels on +// 8+ rank single-node gfx1250 comms. +TEST(WrapMicrotestIsolated, UpdateCollectiveChannels_ReduceScatterSimpleBand) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveChannels_ReduceScatterSimpleBand", + []() { + ncclComm* comm = MakeTuningComm(16); + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 8 * kMiB)); + EXPECT_EQ(128, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 8 * kMiB + 16)); + EXPECT_EQ(128, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 32 * kMiB)); + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 32 * kMiB + 16)); + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_LL128, 16 * kMiB)) << "LL128"; + EXPECT_EQ(192, CappedChannels(comm, ncclFuncAllGather, NCCL_PROTO_SIMPLE, 16 * kMiB)) << "AllGather"; + comm->nNodes = 2; + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 16 * kMiB)) << "2 nodes"; + comm->nNodes = 1; + comm->nRanks = 4; + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 16 * kMiB)) << "4 ranks"; + DeleteCommWithArch(comm); + comm = MakeTuningComm(16, "gfx950"); + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 16 * kMiB)) << "gfx950"; + DeleteCommWithArch(comm); + }); +} + +// A user channel count (NCCL_MAX_NCHANNELS or NCCL_MIN_NCHANNELS) turns the cap off. +TEST(WrapMicrotestIsolated, UpdateCollectiveChannels_UserChannelCountTurnsItOff) { + RUN_ISOLATED_TEST( + "Wrap_UpdateCollectiveChannels_UserChannelCountTurnsItOff", + []() { + ncclComm* comm = MakeTuningComm(16); + g_paramMaxNchannels = 192; + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 16 * kMiB)); + g_paramMaxNchannels = -2; + g_paramMinNchannels = 64; + EXPECT_EQ(192, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 16 * kMiB)); + g_paramMinNchannels = -2; + EXPECT_EQ(128, CappedChannels(comm, ncclFuncReduceScatter, NCCL_PROTO_SIMPLE, 16 * kMiB)) << "control"; + DeleteCommWithArch(comm); + }); +} + // CE registered via symmetric windows (branch #3): distinct from the // force-scratch test above by leaving ceScratch/ddaScratch at their // defaults (unset) so branch #2 can't fire, and instead making ceAvailable diff --git a/projects/rccl/test/test_categories_fixtures_debug.yaml b/projects/rccl/test/test_categories_fixtures_debug.yaml index db2422d55b5f..75ba34f1c596 100644 --- a/projects/rccl/test/test_categories_fixtures_debug.yaml +++ b/projects/rccl/test/test_categories_fixtures_debug.yaml @@ -29,6 +29,7 @@ test_categories: - "DdaAlltoAllThreshold.*" - "DdaIpcEligibilityTest.*" - "DdaFabricEligibilityTest.*" + - "DdaFabricLL128EligibilityTest.*" - "DdaFabricMaxBlocksTest.*" - "DdaFabricScratchSizingTest.*" - "RcclDdaTierThresholds.*" diff --git a/projects/rccl/test/test_categories_micro.yaml b/projects/rccl/test/test_categories_micro.yaml index ae0396c0621b..b1ed2a337b65 100644 --- a/projects/rccl/test/test_categories_micro.yaml +++ b/projects/rccl/test/test_categories_micro.yaml @@ -93,6 +93,7 @@ test_categories: - "GdrProbeTest.*" # ce_coll.cc (rccl-UnitTestsMicroCe) - "CeCollMicrotest.*" + - "LaunchPaths/CeAlltoAllvSyncMicrotest.*" # dev_runtime.cc - "Alloc*" - "Comm*" diff --git a/projects/rccl/tools/topo_expl/Makefile b/projects/rccl/tools/topo_expl/Makefile index da9372c0cddb..ba6cac731282 100644 --- a/projects/rccl/tools/topo_expl/Makefile +++ b/projects/rccl/tools/topo_expl/Makefile @@ -122,6 +122,7 @@ RCCL_FILES = hipify_rccl/graph/topo.cc \ ../../src/misc/param.cc \ ../../src/misc/nvmlwrap_stub.cc \ hipify_rccl/rccl_wrap.cc \ + hipify_rccl/rccl_arch_thresholds.cc \ hipify_rccl/collectives.cc \ hipify_rccl/register.cc \ hipify_rccl/enqueue.cc @@ -161,6 +162,7 @@ hipify: cp -a ../../src/register/register.cc hipify_rccl/ cp -a ../../src/collectives.cc hipify_rccl/ cp -a ../../src/rccl_wrap.cc hipify_rccl/ + cp -a ../../src/rccl_arch_thresholds.cc hipify_rccl/ cp -a ../../src/misc/archinfo.cc hipify_rccl/graph/ # Hipify all header and source files (CUDA -> HIP). # Walk every subdirectory under hipify_rccl/include/ recursively so that