From f2fd1a45b01f6d215b6cdd56ed1d069309cdaf3c Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Thu, 8 Oct 2026 11:25:49 -0500 Subject: [PATCH 1/7] optimize gin allreduce for MI4XX --- projects/rccl/cmake/DeviceLinker.cmake | 113 ++++--- .../rccl/src/algorithms/gin/gin_all_reduce.h | 21 +- .../algorithms/gin/gin_all_reduce_policy.h | 170 ++++++++++- .../algorithms/gin/sdma/all_reduce_gin_sdma.h | 282 ++++++++++++++++++ .../gin/sdma/gin_all_reduce_sdma.cu | 125 +++++++- projects/rccl/src/rccl_wrap.cc | 7 +- .../test/common/GinAllReduceTestHelpers.hpp | 6 +- .../test/gin/gin_all_reduce_policy_test.cpp | 169 +++++++++++ 8 files changed, 802 insertions(+), 91 deletions(-) diff --git a/projects/rccl/cmake/DeviceLinker.cmake b/projects/rccl/cmake/DeviceLinker.cmake index bfb53916a40b..78acd876228d 100644 --- a/projects/rccl/cmake/DeviceLinker.cmake +++ b/projects/rccl/cmake/DeviceLinker.cmake @@ -915,12 +915,63 @@ if(ENABLE_ROCSHMEM_GIN) ) endif() +# =========================================================================== +# rocshmem QP device bitcode for GIN GDA puts (QueuePair::put_nbi, etc.). +# Used by GIN symmetric kernels. The installed per-arch .bc files are +# arch-optimized and cannot be used with -mlink-builtin-bitcode across archs, +# so this is unoptimized IR + constmem stubs. +# =========================================================================== +set(GIN_ROCSHMEM_QP_BC_FLAG "") +set(GIN_ROCSHMEM_QP_DEPS "") +if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR AND DL_GPU_TARGETS) + list(GET DL_GPU_TARGETS 0 _bc_arch) + set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc") + find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED) + find_program(_llvm_dis llvm-dis HINTS ${ROCM_PATH}/llvm/bin REQUIRED) + find_program(_llvm_as llvm-as HINTS ${ROCM_PATH}/llvm/bin REQUIRED) + + set(_cm_src "${CMAKE_SOURCE_DIR}/src/gin/gin_rocshmem_constmem.hip") + set(_cm_bc "${DEVICE_BUILD_DIR}/gin_rocshmem_constmem.bc") + set(_qp_raw "${DEVICE_BUILD_DIR}/rocshmem_qp_raw.bc") + + add_custom_command( + OUTPUT ${_cm_bc} + COMMAND ${DL_CLANG} + -x hip --cuda-device-only --offload-arch=${_bc_arch} + -emit-llvm -Xclang -disable-llvm-passes + -std=c++17 -fPIC + -I${ROCSHMEM_SOURCE_DIR}/src + -I${ROCSHMEM_SOURCE_DIR}/include + -c -o ${_cm_bc} ${_cm_src} + DEPENDS ${_cm_src} + COMMENT "DL: compiling rocshmem constmem stubs to device bitcode" + VERBATIM) + + add_custom_command( + OUTPUT ${_qp_bc} + COMMAND ${_llvm_link} ${_cm_bc} -o ${_qp_raw} + COMMAND ${_llvm_dis} -o ${_qp_raw}.ll ${_qp_raw} + COMMAND grep -v -E "@llvm[.]compiler[.]used|@__hip_cuid_" + ${_qp_raw}.ll > ${_qp_raw}.clean.ll + COMMAND ${_llvm_as} ${_qp_raw}.clean.ll -o ${_qp_bc} + DEPENDS ${_cm_bc} + COMMENT "DL: linking rocshmem QP device bitcode (with constmem, stripped)" + VERBATIM) + add_custom_target(dl_rocshmem_qp_bc DEPENDS ${_qp_bc}) + if(TARGET rocshmem_static) + add_dependencies(dl_rocshmem_qp_bc rocshmem_static) + endif() + set(GIN_ROCSHMEM_QP_BC_FLAG -Xclang -mlink-builtin-bitcode -Xclang ${_qp_bc}) + set(GIN_ROCSHMEM_QP_DEPS dl_rocshmem_qp_bc) +endif() + # =========================================================================== # gin_all_reduce_sdma.cu.cpp: GIN-SDMA allreduce kernel. # # The defines go here because this file is filtered out of the rccl target, so # per-source properties never apply. Leaving only SDMA on keeps ncclGinCallImpl # on its single backend branch. GDA must stay off, as this object gets no QP bitcode. +# gfx1250 GinScatter and gfx950 two-shot both use Anvil SDMA gin.put(). # =========================================================================== set(GIN_ALLREDUCE_SDMA_FAT_OBJ "") if(ENABLE_ROCSHMEM_GIN) @@ -1300,65 +1351,9 @@ endforeach() # =========================================================================== set(SYM_FAT_OBJS "") if(GENERATE_SYM_KERNELS) - # When ENABLE_ROCSHMEM_GIN is set, GIN device templates reference rocshmem - # device symbols (QueuePair::put_nbi, atomic_add, etc.). The installed per-arch - # .bc files are arch-optimized (opt -mcpu=) and can't be used with - # -mlink-builtin-bitcode across archs. Instead, llvm-link the pre-opt - # individual source .bc files (arch-agnostic) into a minimal QP-only bitcode. - set(_sym_rocshmem_bc_flag "") - set(_sym_rocshmem_deps "") - if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR) - # Pick the first arch's pre-opt bitcode dir (all archs produce identical - # unoptimized IR since -Xclang -disable-llvm-passes is used). - list(GET DL_GPU_TARGETS 0 _bc_arch) - set(_bc_dir "${ROCSHMEM_SOURCE_DIR}/build/bitcode/${_bc_arch}") - set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc") - find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED) - find_program(_llvm_dis llvm-dis HINTS ${ROCM_PATH}/llvm/bin REQUIRED) - find_program(_llvm_as llvm-as HINTS ${ROCM_PATH}/llvm/bin REQUIRED) - - # Pipeline: - # 1. Compile gin_rocshmem_constmem.hip → device-only .bc (provides - # rocshmem::constmem and rocshmem::logd_constants definitions that - # queue_pair.bc references as external) - # 2. llvm-link QP .bc files + constmem .bc into one module - # 3. Strip @llvm.compiler.used and @__hip_cuid_ via text round-trip - # (these AMDGCN addrspace(1) appending globals clash with the - # host-side addrspace(0) equivalents in fat-object compilation) - set(_cm_src "${CMAKE_SOURCE_DIR}/src/gin/gin_rocshmem_constmem.hip") - set(_cm_bc "${DEVICE_BUILD_DIR}/gin_rocshmem_constmem.bc") - set(_qp_raw "${DEVICE_BUILD_DIR}/rocshmem_qp_raw.bc") - - add_custom_command( - OUTPUT ${_cm_bc} - COMMAND ${DL_CLANG} - -x hip --cuda-device-only --offload-arch=${_bc_arch} - -emit-llvm -Xclang -disable-llvm-passes - -std=c++17 -fPIC - -I${ROCSHMEM_SOURCE_DIR}/src - -I${ROCSHMEM_SOURCE_DIR}/include - -c -o ${_cm_bc} ${_cm_src} - DEPENDS ${_cm_src} - COMMENT "DL: compiling rocshmem constmem stubs to device bitcode" - VERBATIM) - - add_custom_command( - OUTPUT ${_qp_bc} - COMMAND ${_llvm_link} ${_cm_bc} -o ${_qp_raw} - COMMAND ${_llvm_dis} -o ${_qp_raw}.ll ${_qp_raw} - COMMAND grep -v -E "@llvm[.]compiler[.]used|@__hip_cuid_" - ${_qp_raw}.ll > ${_qp_raw}.clean.ll - COMMAND ${_llvm_as} ${_qp_raw}.clean.ll -o ${_qp_bc} - DEPENDS ${_cm_bc} - COMMENT "DL: linking rocshmem QP device bitcode (with constmem, stripped)" - VERBATIM) - add_custom_target(dl_rocshmem_qp_bc DEPENDS ${_qp_bc}) - if(TARGET rocshmem_static) - add_dependencies(dl_rocshmem_qp_bc rocshmem_static) - endif() - set(_sym_rocshmem_bc_flag -Xclang -mlink-builtin-bitcode -Xclang ${_qp_bc}) - set(_sym_rocshmem_deps dl_rocshmem_qp_bc) - endif() + # GIN symmetric kernels reuse the QP bitcode generated above. + set(_sym_rocshmem_bc_flag "${GIN_ROCSHMEM_QP_BC_FLAG}") + set(_sym_rocshmem_deps "${GIN_ROCSHMEM_QP_DEPS}") file(GLOB _sym_srcs CONFIGURE_DEPENDS "${HIPIFY_DIR}/gensrc/symmetric/*.cpp") foreach(_sym_src IN LISTS _sym_srcs) get_filename_component(_sym_name "${_sym_src}" NAME_WE) diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce.h b/projects/rccl/src/algorithms/gin/gin_all_reduce.h index 53152ed74158..f261cdbf35d6 100755 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce.h @@ -2,10 +2,11 @@ * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. * * Host entry points for the GIN-SDMA AllReduce path selected by rcclSelectAllReduce - * when symmetric windows are used. By default only messages >= 256 MiB - * (GIN two-shot) take this path; smaller messages use DDA AllReduce. + * when symmetric windows are used. By default gfx950 takes GIN two-shot for + * messages >= 256 MiB; gfx1250 with Anvil SDMA takes GinScatter from 64 MiB. + * Smaller messages use DDA AllReduce. * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables LSA one-shot (<= 4 MiB) and - * LSA two-shot ((4 MiB, 256 MiB)). + * LSA two-shot ((4 MiB, GIN floor)). * See LICENSE.txt for license information. ******************************************************************************/ @@ -19,12 +20,14 @@ #include struct ncclComm; +struct ncclDevrWindow; // Size and CTA constants live in gin_all_reduce_policy.h so host unit tests and // ncclAllReduceGinSdmaEligible() share one definition. // LSA one-shot for messages <= kGinAllReduceLsaOneShotMaxBytes (force-enable only). -// LSA two-shot for (4 MiB, 256 MiB) (force-enable only). -// GIN two-shot for messages >= kGinAllReduceGinTwoShotMinBytes (default path). +// LSA two-shot between 4 MiB and the GIN floor (force-enable only). +// gfx950 + Anvil SDMA: GIN two-shot for messages >= kGinAllReduceGinTwoShotMinBytes (256 MiB). +// gfx1250 + Anvil SDMA: GinScatter for messages >= kGinAllReduceGinScatterMinBytes (64 MiB). // Lazily created on the first eligible AllReduce and torn down with the comm. // Declared unconditionally: ncclComm embeds this even when ENABLE_ROCSHMEM_GIN is off. @@ -36,10 +39,16 @@ struct ncclComm; // would freeze at its captured value while the device counters kept advancing. // intraGpuCtaBar is two device uint32s (arrived, sense) for the GIN two-shot intra-GPU // CTA barrier; the pointer is stable, the words themselves advance on the device. +// scratch / scratchWin are the symmetric staging buffer for GinScatter (512 MiB of +// CE-style ping-pong incoming slots). Allocated and window-registered once on +// gfx1250 before GIN connect. struct ncclGinAllReduceState { bool initialized; struct ncclDevComm devComm; uint32_t* intraGpuCtaBar; + void* scratch; + size_t scratchBytes; + struct ncclDevrWindow* scratchWin; }; #if defined(ENABLE_ROCSHMEM_GIN) @@ -51,7 +60,7 @@ bool ncclAllReduceGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* re ncclDataType_t datatype, ncclRedOp_t op); // True when this AllReduce is a GIN-SDMA candidate that the default size policy -// left for DDA (message < 256 MiB and RCCL_GIN_ALLREDUCE_FORCE_ENABLE is not 1). +// left for DDA (message below the GIN floor and RCCL_GIN_ALLREDUCE_FORCE_ENABLE is not 1). bool ncclAllReduceGinSdmaYieldToDda(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op); diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h index 8daf4d49541d..bc576ce6a3a3 100644 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h @@ -1,9 +1,15 @@ /************************************************************************* * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. * - * Size and alignment policy for GIN-SDMA AllReduce. Pure host header included by - * gin_all_reduce.h so unit tests exercise the same gates - * ncclAllReduceGinSdmaEligible() uses after the comm / buffer / gfx950 checks pass. + * Size and alignment policy for GIN-SDMA AllReduce. Included by gin_all_reduce.h + * so unit tests exercise the same gates ncclAllReduceGinSdmaEligible() uses after + * the comm / buffer / gfx950 / gfx1250 checks pass. Chunk-plan helpers are + * host+device so the GinScatter kernel splits the same way the tests do. + * + * gfx950 default: GIN two-shot from 256 MiB (LSA RS + GIN AG, Anvil SDMA). + * gfx1250 + Anvil SDMA: GinScatter from 64 MiB. + * GinScatter scratch follows CE AllReduce (ce_reduce_impl.h / ce_coll.h): + * 512 MiB of incoming ping-pong slots, reduced output in the user recvbuff. * See LICENSE.txt for license information. ******************************************************************************/ @@ -22,9 +28,133 @@ constexpr int kGinAllReduceLsaThreadsPerCta = 512; constexpr size_t kGinAllReduceLsaOneShotMaxBytes = 4ULL * 1024 * 1024; constexpr size_t kGinAllReduceLsaTwoShotMidBytes = 32ULL * 1024 * 1024; constexpr size_t kGinAllReduceGinTwoShotMinBytes = 256ULL * 1024 * 1024; +// GinScatter (CE-style local reduce) from this size inclusive on gfx1250. +constexpr size_t kGinAllReduceGinScatterMinBytes = 64ULL * 1024 * 1024; +// Put/reduce overlap starts here. Below this, ginScatter is one full scatter then +// one full reduce (or sequential slot reuse if scratch cannot hold the shard). +constexpr size_t kGinAllReduceGinScatterPipelineMinBytes = 512ULL * 1024 * 1024; +// Largest AllReduce ginScatter is sized for (2 GiB and 4 GiB included). +constexpr size_t kGinAllReduceGinScatterMaxBytes = 4ULL * 1024 * 1024 * 1024; +// CE ceARTmpBuf analogue: 512 MiB of incoming ping-pong. Reduced shard is recvbuff. +constexpr size_t kGinAllReduceTwoShotScratchBytes = 512ULL * 1024 * 1024; constexpr size_t kGinAllReduceMinPutBytes = 128; +// Pipeline granule when the AllReduce is >= 512 MiB and scratch holds the shard. +// With the default 512 MiB scratch, messages above 512 MiB are staged and the +// slot (scratch / (nSlots * nRanks)) is the granule instead. +constexpr size_t kGinAllReduceGinScatterChunkBytes = 64ULL * 1024 * 1024; +constexpr size_t kGinAllReduceGinScatterChunkAlign = 16; +// Matches NCCL_CE_NUM_SLOTS: ping-pong scatter staging when the shard does not +// fit in one incoming column. +constexpr int kGinAllReduceGinScatterNumSlots = 2; + +#if defined(__HIPCC__) || defined(__CUDACC__) +#define GIN_AR_HD __host__ __device__ __forceinline__ +#else +#define GIN_AR_HD inline +#endif + +GIN_AR_HD size_t ginAllReduceGinScatterAlignDown(size_t n, size_t align) { + if (align == 0) return n; + return (n / align) * align; +} + +// CE ceARTmpBuf: nSlots * nRanks * slotBytes. Reduced output is the user recvbuff. +GIN_AR_HD size_t ginAllReduceGinScatterStagedScratchBytes(size_t slotBytes, int nRanks, int nSlots) { + if (nRanks <= 0 || nSlots <= 0) return 0; + return static_cast(nSlots) * static_cast(nRanks) * slotBytes; +} + +// Largest 16-byte-aligned per-rank slot that fits in scratch with nSlots windows. +GIN_AR_HD size_t ginAllReduceGinScatterFitSlot(size_t scratchBytes, int nRanks, int nSlots) { + if (nRanks <= 0 || nSlots <= 0) return 0; + const size_t denom = static_cast(nSlots) * static_cast(nRanks); + return ginAllReduceGinScatterAlignDown(scratchBytes / denom, kGinAllReduceGinScatterChunkAlign); +} + +// True when scratch cannot hold nRanks full incoming columns (CE pipelined path). +GIN_AR_HD bool ginAllReduceGinScatterStaged(size_t perRankBytes, int nRanks, size_t scratchBytes) { + if (nRanks <= 0) return false; + return scratchBytes < static_cast(nRanks) * perRankBytes; +} + +// Split one per-rank slice. Full-slot path: one chunk below 512 MiB, 16 MiB +// stages at 512 MiB+. Staged path (message > scratch): CE ping-pong — two slots +// and overlap at >= 512 MiB; one slot and sequential reuse below 512 MiB. +GIN_AR_HD void ginAllReduceGinScatterChunkPlan(size_t perRankBytes, int nRanks, size_t scratchBytes, int* nChunks, + size_t* uniformBytes, size_t* lastBytes, int* nPhases) { + const size_t maxChunk = kGinAllReduceGinScatterChunkBytes; + const size_t align = kGinAllReduceGinScatterChunkAlign; + const size_t totalBytes = nRanks > 0 ? perRankBytes * static_cast(nRanks) : 0; + if (perRankBytes == 0) { + *nChunks = 1; + *uniformBytes = 0; + *lastBytes = 0; + *nPhases = 1; + return; + } + if (!ginAllReduceGinScatterStaged(perRankBytes, nRanks, scratchBytes)) { + *nPhases = 1; + if (totalBytes < kGinAllReduceGinScatterPipelineMinBytes || perRankBytes <= maxChunk) { + *nChunks = 1; + *uniformBytes = perRankBytes; + *lastBytes = perRankBytes; + return; + } + int n = static_cast((perRankBytes + maxChunk - 1) / maxChunk); + size_t uniform = (perRankBytes / static_cast(n) / align) * align; + if (uniform < align) uniform = align; + *nChunks = n; + *uniformBytes = uniform; + *lastBytes = perRankBytes - uniform * static_cast(n - 1); + return; + } + + int phases = (totalBytes >= kGinAllReduceGinScatterPipelineMinBytes) ? kGinAllReduceGinScatterNumSlots : 1; + size_t slot = ginAllReduceGinScatterFitSlot(scratchBytes, nRanks, phases); + if (slot < align && phases > 1) { + phases = 1; + slot = ginAllReduceGinScatterFitSlot(scratchBytes, nRanks, 1); + } + if (slot < align) { + *nChunks = 1; + *uniformBytes = perRankBytes; + *lastBytes = perRankBytes; + *nPhases = 1; + return; + } + if (perRankBytes <= slot) { + *nChunks = 1; + *uniformBytes = perRankBytes; + *lastBytes = perRankBytes; + *nPhases = 1; + return; + } + int n = static_cast((perRankBytes + slot - 1) / slot); + *nChunks = n; + *uniformBytes = slot; + *lastBytes = perRankBytes - slot * static_cast(n - 1); + *nPhases = phases; +} + +GIN_AR_HD size_t ginAllReduceGinScatterChunkOff(int chunk, size_t uniformBytes) { + return static_cast(chunk) * uniformBytes; +} + +GIN_AR_HD size_t ginAllReduceGinScatterChunkSize(int chunk, int nChunks, size_t uniformBytes, size_t lastBytes) { + return (chunk == nChunks - 1) ? lastBytes : uniformBytes; +} + +#undef GIN_AR_HD + +// Full-slot incoming columns (no reduced scratch column; CE writes the reduced +// shard to recvbuff). Used to decide whether the message fits without staging. +inline size_t ginAllReduceGinScatterScratchBytes(size_t chunkBytes, int nRanks) { + if (nRanks <= 0) return 0; + return static_cast(nRanks) * chunkBytes; +} + // Two-shot kernels require a whole number of elements per rank and a 16-byte // per-rank slice so vector loads stay aligned. inline bool ginAllReduceTwoShotEligible(size_t count, size_t typeSize, int nRanks) { @@ -40,26 +170,44 @@ inline bool ginAllReduceGinTwoShotEligible(size_t count, size_t typeSize, int nR return chunkBytes >= kGinAllReduceMinPutBytes; } +// ginScatter from 64 MiB when scratch holds either nRanks full incoming columns +// or at least one 16-byte-aligned staging slot per rank. +inline bool ginAllReduceGinScatterLaunch(size_t bytes, size_t chunkBytes, int nRanks, size_t scratchBytes) { + if (bytes < kGinAllReduceGinScatterMinBytes) return false; + if (nRanks <= 0) return false; + if (scratchBytes >= ginAllReduceGinScatterScratchBytes(chunkBytes, nRanks)) return true; + return ginAllReduceGinScatterFitSlot(scratchBytes, nRanks, 1) >= kGinAllReduceGinScatterChunkAlign; +} + // Size-only GIN AllReduce policy. forceEnable is RCCL_GIN_ALLREDUCE_FORCE_ENABLE==1. -// Default: true only for messages >= 256 MiB that pass GIN two-shot alignment. -// Force: LSA one-shot <= 4 MiB (and >= 512 KiB), LSA two-shot in between, GIN two-shot at 256 MiB+. -inline bool ginAllReduceSizePolicyEligible(size_t count, size_t typeSize, int nRanks, bool forceEnable) { +// ginMinBytes is the default-path floor (256 MiB two-shot on gfx950, 64 MiB +// GinScatter on gfx1250). Force: LSA one-shot <= 4 MiB (and >= 512 KiB), LSA +// two-shot in between, GIN path at ginMinBytes+. +inline bool ginAllReduceSizePolicyEligible(size_t count, size_t typeSize, int nRanks, bool forceEnable, + size_t ginMinBytes = kGinAllReduceGinTwoShotMinBytes) { const size_t bytes = count * typeSize; - if (bytes < kGinAllReduceGinTwoShotMinBytes && !forceEnable) return false; + if (bytes < ginMinBytes && !forceEnable) return false; if (bytes < static_cast(kGinAllReduceMinBytes)) return false; // Inclusive 4 MiB: one-shot has no per-rank alignment requirement. Two-shot // starts strictly above this, matching ncclAllReduceGinSdmaTyped(). if (bytes <= kGinAllReduceLsaOneShotMaxBytes) return true; - if (bytes >= kGinAllReduceGinTwoShotMinBytes) { + if (bytes >= ginMinBytes) { return ginAllReduceGinTwoShotEligible(count, typeSize, nRanks); } return ginAllReduceTwoShotEligible(count, typeSize, nRanks); } -// Smaller-than-256 MiB GIN candidates yield to DDA unless FORCE_ENABLE=1. -inline bool ginAllReduceYieldToDdaBySize(size_t count, size_t typeSize, bool forceEnable) { +// Smaller-than-ginMinBytes GIN candidates yield to DDA unless FORCE_ENABLE=1. +inline bool ginAllReduceYieldToDdaBySize(size_t count, size_t typeSize, bool forceEnable, + size_t ginMinBytes = kGinAllReduceGinTwoShotMinBytes) { if (forceEnable) return false; - return count * typeSize < kGinAllReduceGinTwoShotMinBytes; + return count * typeSize < ginMinBytes; +} + +// Host dispatch for ginAllReduceTwoShotGinScatterKernel: gfx1250 and >= 64 MiB. +// Eligibility still requires the Anvil SDMA GIN backend before launch. +inline bool ginAllReduceGinScatterDispatch(bool isGfx1250, size_t bytes) { + return isGfx1250 && bytes >= kGinAllReduceGinScatterMinBytes; } #endif diff --git a/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h b/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h index 62c786f09bc6..2be0cad01342 100755 --- a/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h +++ b/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h @@ -5,6 +5,11 @@ * allReduceLsaOneShotKernel — read all peers, sum, write all peers (small messages). * lsaAllReduceTwoShotKernel — LSA reduce-scatter, barrier, LSA all-gather (remote reads only). * ginAllReduceTwoShotKernel — LSA reduce-scatter, intra-GPU CTA barrier, GIN all-gather from recv. + * ginAllReduceTwoShotGinScatterKernel — CTA 0 GIN-puts remote send slices + * (self from send, remotes from scratch). Scratch is 512 MiB of CE-style + * ping-pong incoming slots; reduced output is recvbuff. Local reduce matches + * CE (16B vectors, GpuUnroll=4, ranks 0..nRanks-1). Used on gfx1250 from 64 MiB + * when the Anvil SDMA GIN backend is enabled. * * One-shot follows projects/rccl-tests/src/all_reduce.cu allReduceLsaKernel. ******************************************************************************/ @@ -14,6 +19,7 @@ #include "nccl_device.h" #include "nccl_device/ptr.h" #include "algorithms/dda/device/CollCommon.h" +#include "algorithms/gin/gin_all_reduce_policy.h" namespace gin::sdma { using dda::common::vecElementAdd; @@ -330,4 +336,280 @@ __launch_bounds__(512) } +// GIN-put one pipeline chunk of send[dst] into dest dst's staging slot for every dst != rank. +__device__ __forceinline__ void ginScatterPutRemoteChunk(ncclGin& gin, ncclTeam world, int rank, int nRanks, + ncclWindow_t scratchWin, size_t destSlotOff, + ncclWindow_t sendWin, size_t sendOff, size_t perRankBytes, + size_t sendChunkOff, size_t chunkBytes, unsigned signalIndex) { + for (int dst = static_cast(threadIdx.x); dst < nRanks; dst += static_cast(blockDim.x)) { + if (dst == rank) { + continue; + } + gin.put(world, dst, scratchWin, destSlotOff, sendWin, + sendOff + static_cast(dst) * perRankBytes + sendChunkOff, chunkBytes, + ncclGin_WeakSignalInc{signalIndex}); + } +} + +__device__ __forceinline__ size_t ginScatterPutDestOff(size_t scratchOff, int rank, int nRanks, size_t perRankBytes, + size_t slotBytes, int slot, size_t chunkOff, bool staged) { + if (!staged) { + return scratchOff + static_cast(rank) * perRankBytes + chunkOff; + } + return scratchOff + static_cast(slot) * static_cast(nRanks) * slotBytes + + static_cast(rank) * slotBytes; +} + +__device__ __forceinline__ const char* ginScatterIncomingBase(const char* scratchBase, int nRanks, size_t slotBytes, + int slot, bool staged) { + if (!staged) { + return scratchBase; + } + return scratchBase + static_cast(slot) * static_cast(nRanks) * slotBytes; +} + +// CE ncclCeLocalReduceKernelVec: 16B vectors, GpuUnroll=4, rank loop 0..nRanks-1 +// into recvbuff + rank*shard. Self is not written to scratch (GIN-put skips it), +// so rank == myRank loads from send instead of the incoming slot. +template +__device__ __forceinline__ const char* ginScatterReduceSrc(int peer, int rank, const char* selfSend, + const char* incomingBase, size_t incomingStride, size_t elem, + size_t elemStart, bool staged) { + if (peer == rank) { + return selfSend + elem * sizeof(T); + } + const size_t srcOff = staged ? (elem - elemStart) * sizeof(T) : elem * sizeof(T); + return incomingBase + static_cast(peer) * incomingStride + srcOff; +} + +template +__device__ __forceinline__ void ginScatterReduceChunk(const char* selfSend, const char* incomingBase, char* reducedOut, + int rank, int nRanks, size_t incomingStride, size_t elemStart, + size_t elemEnd, bool staged) { + constexpr int W = static_cast(sizeof(uint4) / sizeof(T)); + constexpr int U = 4; + const size_t nVec = (elemEnd - elemStart) / static_cast(W); + const size_t tid = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + const size_t stride = static_cast(blockDim.x) * gridDim.x; + + size_t vi = tid; + for (; vi + static_cast(U - 1) * stride < nVec; vi += stride * static_cast(U)) { + uint4 acc[U]; + size_t elem[U]; +#pragma unroll + for (int i = 0; i < U; ++i) { + elem[i] = elemStart + (vi + static_cast(i) * stride) * static_cast(W); + acc[i] = lsaLoadVec(ginScatterReduceSrc(0, rank, selfSend, incomingBase, incomingStride, elem[i], elemStart, + staged)); + } +#pragma clang loop unroll(disable) vectorize(disable) + for (int r = 1; r < nRanks; ++r) { + uint4 tmp[U]; +#pragma unroll + for (int i = 0; i < U; ++i) { + tmp[i] = lsaLoadVec(ginScatterReduceSrc(r, rank, selfSend, incomingBase, incomingStride, elem[i], elemStart, + staged)); + } +#pragma unroll + for (int i = 0; i < U; ++i) { + acc[i] = vecElementAdd(acc[i], tmp[i]); + } + } +#pragma unroll + for (int i = 0; i < U; ++i) { + lsaStoreVec(reducedOut + elem[i] * sizeof(T), acc[i]); + } + } + + const size_t vectorTailStart = nVec - (nVec % (stride * static_cast(U))); + if (nVec > 0 && nVec % (stride * static_cast(U)) != 0) { + for (size_t vIdx = vectorTailStart + tid; vIdx < nVec; vIdx += stride) { + const size_t e = elemStart + vIdx * static_cast(W); + uint4 acc = + lsaLoadVec(ginScatterReduceSrc(0, rank, selfSend, incomingBase, incomingStride, e, elemStart, staged)); +#pragma clang loop unroll(disable) vectorize(disable) + for (int r = 1; r < nRanks; ++r) { + acc = vecElementAdd( + acc, lsaLoadVec(ginScatterReduceSrc(r, rank, selfSend, incomingBase, incomingStride, e, elemStart, + staged))); + } + lsaStoreVec(reducedOut + e * sizeof(T), acc); + } + } +} + +__device__ __forceinline__ void ginScatterIssuePut(ncclGin& gin, ncclTeam world, int rank, int nRanks, + ncclWindow_t scratchWin, size_t scratchOff, ncclWindow_t sendWin, + size_t sendOff, size_t perRankBytes, size_t slotBytes, + size_t uniformBytes, size_t lastBytes, int nChunks, int nPhases, + bool staged, int chunk, unsigned signalIndex) { + const size_t chunkOff = ginAllReduceGinScatterChunkOff(chunk, uniformBytes); + const size_t chunkBytes = ginAllReduceGinScatterChunkSize(chunk, nChunks, uniformBytes, lastBytes); + const int slot = staged ? ((nPhases <= 1) ? 0 : (chunk % nPhases)) : 0; + const size_t destOff = + ginScatterPutDestOff(scratchOff, rank, nRanks, perRankBytes, slotBytes, slot, staged ? 0 : chunkOff, staged); + ginScatterPutRemoteChunk(gin, world, rank, nRanks, scratchWin, destOff, sendWin, sendOff, perRankBytes, chunkOff, + chunkBytes, signalIndex); +} + +// Thread 0 polls; callers that need the rest of the CTA to observe completion +// follow with ginIntraGpuCtaBarrier or __syncthreads. +__device__ __forceinline__ void ginScatterWaitThread0(ncclGin& gin, unsigned signalIndex, uint64_t least) { + if (threadIdx.x == 0) { + gin.waitSignal(ncclCoopThread(), signalIndex, least); + } +} + +// One-way doorbell: this rank finished reading a reused scratch slot. Cheaper than +// a world barrier (CE uses per-slot signal==0 the same way). +__device__ __forceinline__ void ginScatterNotifySlotFree(ncclGin& gin, ncclTeam world, int rank, int nRanks, + unsigned doneSignal) { + for (int dst = static_cast(threadIdx.x); dst < nRanks; dst += static_cast(blockDim.x)) { + if (dst == rank) { + continue; + } + gin.signal(world, dst, ncclGin_WeakSignalInc{doneSignal}); + } +} + +// Two-shot AllReduce over GIN-SDMA. Grid is 56 CTAs x 512 threads. +// +// Scratch follows CE AllReduce (512 MiB, NCCL_CE_NUM_SLOTS ping-pong): +// Fit: [peer 0 .. nRanks) full incoming columns (reduced shard is recvbuff) +// Staged: [slot 0: nRanks slots][slot 1: nRanks slots], slot = step % nPhases +// Default 512 MiB scratch stages ginScatter messages above 512 MiB. +// +// Shot 1: CTA 0 GIN-puts send[dst] into dest's staging slot (dst != myRank). +// Local reduce is CE: 16B vectors, unroll 4, rank 0..nRanks-1 into recv[myRank] +// (self from send). Below 512 MiB a step completes scatter before reduce. At +// 512 MiB+ put of chunk i+1 overlaps reduce of chunk i. CTA 0 GIN all-gathers. +// +// Sync: one world acquire so all ranks snapshot GIN signal baselines. Each pipeline +// step is waitSignal on fused WeakSignalInc (scatter/AG) or a one-way slot-free +// doorbell — not ginBar.sync or gin.flush. Intra-GPU CTA bars only join the 56 CTAs. +template +#if defined(USE_ROCM) +__launch_bounds__(512) +#endif + __global__ void ginAllReduceTwoShotGinScatterKernel(struct ncclDevComm devComm, ncclWindow_t sendWin, size_t sendOff, + ncclWindow_t recvWin, size_t recvOff, ncclWindow_t scratchWin, + size_t scratchOff, size_t countPerRank, int nRanks, + size_t scratchBytes, uint32_t* intraGpuCtaBar) { + constexpr int ginContext = 0; + constexpr unsigned scatterSignal = 0; + constexpr unsigned agSignal = 1; + constexpr unsigned doneSignal = 2; + + // gin/ginBar must outlive the first put; only CTA 0 snapshots signals and world-syncs. + // Pipeline steps use waitSignal / WeakSignalInc (CE doorbells), not world barriers. + ncclGin gin{devComm, ginContext}; + ncclTeam world = ncclTeamWorld(devComm); + ncclCoopCta cta; + ncclBarrierSession ginBar{cta, ncclTeamTagWorld(), gin, blockIdx.x}; + + const size_t perRankBytes = countPerRank * sizeof(T); + const size_t sliceRecvByteOff = recvOff + static_cast(devComm.rank) * perRankBytes; + char* scratchBase = reinterpret_cast(ncclGetLocalPointer(scratchWin, scratchOff)); + char* reducedOut = reinterpret_cast(ncclGetLocalPointer(recvWin, sliceRecvByteOff)); + const char* selfSend = reinterpret_cast(ncclGetLocalPointer(sendWin, sendOff)) + + static_cast(devComm.rank) * perRankBytes; + + int nChunks = 1; + int nPhases = 1; + size_t uniformBytes = perRankBytes; + size_t lastBytes = perRankBytes; + ginAllReduceGinScatterChunkPlan(perRankBytes, nRanks, scratchBytes, &nChunks, &uniformBytes, &lastBytes, &nPhases); + const bool staged = ginAllReduceGinScatterStaged(perRankBytes, nRanks, scratchBytes); + const size_t slotBytes = staged ? ginAllReduceGinScatterFitSlot(scratchBytes, nRanks, nPhases) : perRankBytes; + const size_t incomingStride = staged ? slotBytes : perRankBytes; + const bool overlap = (nChunks > 1) && !(staged && nPhases <= 1); + + const uint64_t signalsPerChunk = static_cast(nRanks - 1); + uint64_t scatterValue = 0; + uint64_t doneValue = 0; + uint64_t agValue = 0; + + // One world acquire so every rank snapshots signal baselines before any put. + // Later pipeline steps are waitSignal doorbells (CE), not ginBar.sync. + if (blockIdx.x == 0) { + scatterValue = gin.readSignal(scatterSignal); + doneValue = gin.readSignal(doneSignal); + agValue = gin.readSignal(agSignal); + ginBar.sync(cta, cuda::memory_order_acquire, ncclGinFenceLevel::None); + } + + if (overlap) { + if (blockIdx.x == 0) { + ginScatterIssuePut(gin, world, devComm.rank, nRanks, scratchWin, scratchOff, sendWin, sendOff, perRankBytes, + slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, 0, scatterSignal); + ginScatterWaitThread0(gin, scatterSignal, scatterValue + signalsPerChunk); + } + ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); + + for (int c = 0; c < nChunks; ++c) { + if (blockIdx.x == 0 && c + 1 < nChunks) { + if (staged && c + 1 >= nPhases) { + ginScatterWaitThread0(gin, doneSignal, doneValue + static_cast(c + 2 - nPhases) * signalsPerChunk); + __syncthreads(); + } + ginScatterIssuePut(gin, world, devComm.rank, nRanks, scratchWin, scratchOff, sendWin, sendOff, perRankBytes, + slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, c + 1, scatterSignal); + } + const size_t chunkOff = ginAllReduceGinScatterChunkOff(c, uniformBytes); + const size_t chunkBytes = ginAllReduceGinScatterChunkSize(c, nChunks, uniformBytes, lastBytes); + const size_t elemStart = chunkOff / sizeof(T); + const size_t elemEnd = elemStart + chunkBytes / sizeof(T); + const int slot = staged ? ((nPhases <= 1) ? 0 : (c % nPhases)) : 0; + ginScatterReduceChunk(selfSend, ginScatterIncomingBase(scratchBase, nRanks, slotBytes, slot, staged), + reducedOut, devComm.rank, nRanks, incomingStride, elemStart, elemEnd, staged); + if (blockIdx.x == 0 && c + 1 < nChunks) { + ginScatterWaitThread0(gin, scatterSignal, scatterValue + static_cast(c + 2) * signalsPerChunk); + } + ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); + if (blockIdx.x == 0 && staged && c + nPhases < nChunks) { + ginScatterNotifySlotFree(gin, world, devComm.rank, nRanks, doneSignal); + __syncthreads(); + } + } + } else { + for (int c = 0; c < nChunks; ++c) { + if (blockIdx.x == 0) { + if (c > 0 && staged) { + ginScatterWaitThread0(gin, doneSignal, doneValue + static_cast(c) * signalsPerChunk); + __syncthreads(); + } + ginScatterIssuePut(gin, world, devComm.rank, nRanks, scratchWin, scratchOff, sendWin, sendOff, perRankBytes, + slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, c, scatterSignal); + ginScatterWaitThread0(gin, scatterSignal, scatterValue + static_cast(c + 1) * signalsPerChunk); + } + ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); + const size_t chunkOff = ginAllReduceGinScatterChunkOff(c, uniformBytes); + const size_t chunkBytes = ginAllReduceGinScatterChunkSize(c, nChunks, uniformBytes, lastBytes); + const size_t elemStart = chunkOff / sizeof(T); + const size_t elemEnd = elemStart + chunkBytes / sizeof(T); + const int slot = staged ? ((nPhases <= 1) ? 0 : (c % nPhases)) : 0; + ginScatterReduceChunk(selfSend, ginScatterIncomingBase(scratchBase, nRanks, slotBytes, slot, staged), + reducedOut, devComm.rank, nRanks, incomingStride, elemStart, elemEnd, staged); + ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); + if (blockIdx.x == 0 && staged && c + 1 < nChunks) { + ginScatterNotifySlotFree(gin, world, devComm.rank, nRanks, doneSignal); + __syncthreads(); + } + } + } + + if (blockIdx.x == 0) { + // Shot 2: waitSignal on fused AG puts is the completion doorbell; no world barrier. + for (int dst = static_cast(threadIdx.x); dst < nRanks; dst += static_cast(blockDim.x)) { + if (dst == devComm.rank) { + continue; + } + gin.put(world, dst, recvWin, sliceRecvByteOff, recvWin, sliceRecvByteOff, perRankBytes, + ncclGin_WeakSignalInc{agSignal}); + } + ginScatterWaitThread0(gin, agSignal, agValue + static_cast(nRanks - 1)); + __syncthreads(); + } +} + } // namespace gin::sdma diff --git a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu index ef95002ef237..2b084921c839 100755 --- a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu +++ b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu @@ -2,12 +2,13 @@ * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. * * GIN-SDMA AllReduce for single-node (scaleup-only) symmetric windows. - * Default: only messages >= 256 MiB take this path (GIN two-shot); smaller - * messages fall through to DDA AllReduce. RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 - * also enables the LSA bands: + * Default: gfx950 messages >= 256 MiB take GIN two-shot (LSA reduce-scatter + + * GIN all-gather). gfx1250 messages >= 64 MiB take GinScatter over Anvil SDMA. + * Smaller messages fall through to DDA AllReduce. + * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables the LSA bands: * <= 4 MiB — LSA one-shot - * (4, 256) MiB — LSA two-shot - * >= 256 MiB — GIN two-shot (LSA reduce-scatter + GIN all-gather) + * (4 MiB, GIN floor) — LSA two-shot + * >= GIN floor — GIN two-shot (gfx950) or GinScatter (gfx1250) * * Compiled with NCCL_GIN_ANVIL_SDMA_ENABLE=1 and NCCL_GIN_PROXY_ENABLE=0 so * ncclGinCallImpl resolves the SDMA backend at compile time. @@ -30,9 +31,14 @@ #include NCCL_PARAM(GinAllReduceEnable, "GIN_ALLREDUCE_ENABLE", 1); -// When 0 (default), GIN AllReduce is eligible only for messages >= 256 MiB. +// When 0 (default), GIN AllReduce is eligible only for messages >= the GIN floor +// (256 MiB two-shot on gfx950, 64 MiB GinScatter on gfx1250). // Set to 1 to also take LSA one-shot / LSA two-shot for smaller sizes. RCCL_PARAM(GinAllReduceForceEnable, "GIN_ALLREDUCE_FORCE_ENABLE", 0); +// Symmetric 512 MiB scratch for GinScatter CE ping-pong incoming slots. +// Allocated once on gfx1250 before GIN connect. 0 skips allocation. +NCCL_PARAM(GinAllReduceScratchBytes, "GIN_ALLREDUCE_SCRATCH_BYTES", + static_cast(kGinAllReduceTwoShotScratchBytes)); // LSA two-shot tuning. CTAs default to kGinAllReduceLsaTwoShotCtasPerPeer * nRanks; the DDA IPC // kernels top out at DDA_IPC_MAXBLOCKS (24), so this range is worth sweeping. NCCL_PARAM(GinAllReduceLsaTwoShotCtas, "GIN_ALLREDUCE_LSA_TWOSHOT_CTAS", 0); @@ -41,6 +47,64 @@ namespace { constexpr bool kSdmaDeviceBackendCompiled = (NCCL_GIN_ANVIL_SDMA_ENABLE != 0); +static bool ginAllReduceIsGfx1250(ncclComm* comm) { + return comm != nullptr && IsArchMatch(comm->archName, "gfx1250"); +} + +static size_t ginAllReduceGinMinBytes(ncclComm* comm) { + return ginAllReduceIsGfx1250(comm) ? kGinAllReduceGinScatterMinBytes : kGinAllReduceGinTwoShotMinBytes; +} + +static bool ginAllReduceScratchWinReady(ncclComm* comm) { + struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + return state->scratchWin != nullptr && state->scratchWin->vidmem != nullptr; +} + +static ncclResult_t ginAllReduceFreeScratch(ncclComm* comm) { + struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + if (state->scratchWin != nullptr) { + NCCLCHECK(ncclCommWindowDeregister(comm, state->scratchWin->vidmem)); + state->scratchWin = nullptr; + } + if (state->scratch != nullptr) { + NCCLCHECK(ncclMemFree(state->scratch)); + state->scratch = nullptr; + } + state->scratchBytes = 0; + return ncclSuccess; +} + +static ncclResult_t ginAllReduceScratchInitOnce(ncclComm* comm) { + struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + if (state->scratch != nullptr) { + return ncclSuccess; + } + + const int64_t requested = ncclParamGinAllReduceScratchBytes(); + if (requested <= 0) { + return ncclSuccess; + } + const size_t allocBytes = static_cast(requested); + + void* scratch = nullptr; + ncclWindow_t scratchWinDev = nullptr; + NCCLCHECK(ncclMemAlloc(&scratch, allocBytes)); + NCCLCHECK(ncclDevrWindowRegisterInGroup(comm, scratch, allocBytes, NCCL_WIN_COLL_SYMMETRIC, &scratchWinDev)); + NCCLCHECK(ncclDevrFindWindow(comm, scratch, &state->scratchWin)); + if (scratchWinDev == nullptr || !ginAllReduceScratchWinReady(comm)) { + if (scratchWinDev != nullptr) { + (void)ncclCommWindowDeregister(comm, scratchWinDev); + } + state->scratchWin = nullptr; + NCCLCHECK(ncclMemFree(scratch)); + return ncclInternalError; + } + state->scratch = scratch; + state->scratchBytes = allocBytes; + INFO(NCCL_INIT, "GIN AllReduce GinScatter scratch: %zu bytes", allocBytes); + return ncclSuccess; +} + // Runs on the first eligible AllReduce, which may itself be inside a graph capture. Everything // here has to stay out of the captured graph and must not disturb the capture in progress: // @@ -57,6 +121,11 @@ constexpr bool kSdmaDeviceBackendCompiled = (NCCL_GIN_ANVIL_SDMA_ENABLE != 0); // barrier before every rank has finished initializing. static ncclResult_t ncclGinAllReduceInitOnce(ncclComm* comm) { NCCLCHECK(ncclDevrInitOnce(comm)); + // Window-register GinScatter scratch while GIN is still off so + // ncclDevrCommCreateInternal GIN-registers every memHead entry, including this one. + if (ginAllReduceIsGfx1250(comm)) { + NCCLCHECK(ginAllReduceScratchInitOnce(comm)); + } struct ncclGinAllReduceState* state = &comm->ginAllReduceState; if (state->initialized) { return ncclSuccess; @@ -200,6 +269,39 @@ static ncclResult_t ncclAllReduceGinSdmaGinTwoShotTyped(const void* sendbuff, vo return ncclSuccess; } +template +static ncclResult_t ncclAllReduceGinSdmaGinScatterTyped(const void* sendbuff, void* recvbuff, size_t count, + ncclComm* comm, cudaStream_t stream, + struct ncclDevrWindow* sendWin, + struct ncclDevrWindow* recvWin) { + NCCLCHECK(ncclGinAllReduceInitOnce(comm)); + + const size_t sendOff = + static_cast(static_cast(sendbuff) - static_cast(sendWin->userPtr)); + const size_t recvOff = + static_cast(static_cast(recvbuff) - static_cast(recvWin->userPtr)); + const size_t countPerRank = count / static_cast(comm->nRanks); + const size_t chunkBytes = countPerRank * sizeof(T); + + struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + if (!ginAllReduceScratchWinReady(comm)) { + WARN("GIN AllReduce GinScatter scratch is not registered"); + return ncclInvalidUsage; + } + if (!ginAllReduceGinScatterLaunch(count * sizeof(T), chunkBytes, comm->nRanks, state->scratchBytes)) { + WARN("GIN AllReduce GinScatter scratch %zu bytes cannot hold message %zu bytes", state->scratchBytes, + count * sizeof(T)); + return ncclInvalidUsage; + } + + gin::sdma::ginAllReduceTwoShotGinScatterKernel + <<>>( + state->devComm, sendWin->vidmem, sendOff, recvWin->vidmem, recvOff, state->scratchWin->vidmem, + /*scratchOff=*/0, countPerRank, comm->nRanks, state->scratchBytes, state->intraGpuCtaBar); + CUDACHECK(cudaGetLastError()); + return ncclSuccess; +} + template static ncclResult_t ncclAllReduceGinSdmaTyped(const void* sendbuff, void* recvbuff, size_t count, ncclComm* comm, cudaStream_t stream, struct ncclDevrWindow* sendWin, @@ -210,6 +312,9 @@ static ncclResult_t ncclAllReduceGinSdmaTyped(const void* sendbuff, void* recvbu if (bytes <= kGinAllReduceLsaOneShotMaxBytes) { return ncclAllReduceGinSdmaOneShotTyped(sendbuff, recvbuff, count, comm, stream, sendWin, recvWin); } + if (ginAllReduceGinScatterDispatch(ginAllReduceIsGfx1250(comm), bytes)) { + return ncclAllReduceGinSdmaGinScatterTyped(sendbuff, recvbuff, count, comm, stream, sendWin, recvWin); + } if (bytes >= kGinAllReduceGinTwoShotMinBytes) { return ncclAllReduceGinSdmaGinTwoShotTyped(sendbuff, recvbuff, count, comm, stream, sendWin, recvWin); } @@ -225,7 +330,7 @@ static bool ginAllReduceBaseEligible(ncclComm* comm, const void* sendbuff, void* if (comm == nullptr || sendbuff == nullptr || recvbuff == nullptr) return false; if (count == 0) return false; - if (!IsArchMatch(comm->archName, "gfx950")) return false; + if (!IsArchMatch(comm->archName, "gfx950") && !ginAllReduceIsGfx1250(comm)) return false; if (op != ncclSum) return false; if (datatype != ncclFloat32 && datatype != ncclFloat16 && datatype != ncclBfloat16) return false; if (!comm->symmetricSupport) return false; @@ -253,13 +358,14 @@ bool ncclAllReduceGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* re ncclDataType_t datatype, ncclRedOp_t op) { if (!ginAllReduceBaseEligible(comm, sendbuff, recvbuff, count, datatype, op)) return false; return ginAllReduceSizePolicyEligible(count, ncclTypeSize(datatype), comm->nRanks, - rcclParamGinAllReduceForceEnable() == 1); + rcclParamGinAllReduceForceEnable() == 1, ginAllReduceGinMinBytes(comm)); } bool ncclAllReduceGinSdmaYieldToDda(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op) { if (!ginAllReduceBaseEligible(comm, sendbuff, recvbuff, count, datatype, op)) return false; - return ginAllReduceYieldToDdaBySize(count, ncclTypeSize(datatype), rcclParamGinAllReduceForceEnable() == 1); + return ginAllReduceYieldToDdaBySize(count, ncclTypeSize(datatype), rcclParamGinAllReduceForceEnable() == 1, + ginAllReduceGinMinBytes(comm)); } ncclResult_t ncclAllReduceGinSdma(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, @@ -284,6 +390,7 @@ ncclResult_t ncclAllReduceGinSdma(const void* sendbuff, void* recvbuff, size_t c ncclResult_t ncclGinAllReduceFinalize(ncclComm* comm) { struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + NCCLCHECK(ginAllReduceFreeScratch(comm)); if (state->intraGpuCtaBar != nullptr) { NCCLCHECK(ncclCudaFree(state->intraGpuCtaBar, comm->memManager)); state->intraGpuCtaBar = nullptr; diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 108cba404082..2f918707dfa4 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -1426,9 +1426,10 @@ ncclResult_t rcclSelectAllReduce(struct ncclComm* comm, const void* sendbuff, vo // registered operands to DDA. // // GIN AllReduce is selected first in this function and requires symmetric - // windows. By default it only claims messages >= 256 MiB, so DDA must still be - // allowed for smaller symmetric AllReduces (otherwise they would hit the - // symmetric kernel instead of DDA). FORCE_ENABLE=1 keeps the original + // windows. By default gfx950 claims messages >= 256 MiB and gfx1250 with + // Anvil SDMA claims >= 64 MiB, so DDA must still be allowed for smaller + // symmetric AllReduces (otherwise they would hit the symmetric kernel + // instead of DDA). FORCE_ENABLE=1 keeps the original // !symEligible gate because GIN already returned above for those sizes. // symSuppressedByMin: DDA wins below symMinR2[AR]; do not block it with symkRequested. bool ddaSymEligible = symkRequested && !symSuppressedByMin; diff --git a/projects/rccl/test/common/GinAllReduceTestHelpers.hpp b/projects/rccl/test/common/GinAllReduceTestHelpers.hpp index 45170074589c..a7c9da352f8a 100644 --- a/projects/rccl/test/common/GinAllReduceTestHelpers.hpp +++ b/projects/rccl/test/common/GinAllReduceTestHelpers.hpp @@ -14,9 +14,9 @@ namespace RcclUnitTesting { // Minimal ncclComm for GIN AllReduce eligibility early-outs that run *before* -// isSymmetricKernelRequested(). Do not set symmetricSupport=true with gfx950, -// ncclSum, and a supported datatype together: that path calls ncclSymkInitOnce -// on this stand-in and is not safe. +// isSymmetricKernelRequested(). Do not set symmetricSupport=true with gfx950 +// or gfx1250, ncclSum, and a supported datatype together: that path calls +// ncclSymkInitOnce on this stand-in and is not safe. struct GinAllReduceMockComm { ncclComm comm{}; diff --git a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp index 3870ea0f6787..ad63c097da92 100644 --- a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp +++ b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp @@ -27,8 +27,14 @@ TEST(GinAllReducePolicy, ThresholdsMatchDocumentedBands) { EXPECT_EQ(kGinAllReduceMinBytes, static_cast(512 * kKiB)); EXPECT_EQ(kGinAllReduceLsaOneShotMaxBytes, 4ull * kMiB); EXPECT_EQ(kGinAllReduceGinTwoShotMinBytes, 256ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterMinBytes, 64ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterPipelineMinBytes, 512ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterMaxBytes, 4ull * 1024ull * kMiB); EXPECT_EQ(kGinAllReduceMaxRanks, 8); EXPECT_EQ(kGinAllReduceMinPutBytes, 128u); + EXPECT_EQ(kGinAllReduceTwoShotScratchBytes, 512ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterChunkBytes, 16ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterNumSlots, 2); } TEST(GinAllReducePolicy, DefaultRejectsBelow256MiB) { @@ -107,4 +113,167 @@ TEST(GinAllReducePolicy, TwoShotRejectsCountNotDivisibleByRanks) { EXPECT_FALSE(ginAllReduceTwoShotEligible(7, kFloat, kRanks)); } +TEST(GinAllReducePolicy, Gfx1250ScatterMinBytesFrom64MiB) { + EXPECT_FALSE( + ginAllReduceSizePolicyEligible(countForBytes(64 * kMiB - 16), kFloat, kRanks, false, kGinAllReduceGinScatterMinBytes)); + EXPECT_TRUE( + ginAllReduceSizePolicyEligible(countForBytes(64 * kMiB), kFloat, kRanks, false, kGinAllReduceGinScatterMinBytes)); + EXPECT_TRUE( + ginAllReduceSizePolicyEligible(countForBytes(128 * kMiB), kFloat, kRanks, false, kGinAllReduceGinScatterMinBytes)); + // gfx950 default floor stays 256 MiB. + EXPECT_FALSE(ginAllReduceSizePolicyEligible(countForBytes(64 * kMiB), kFloat, kRanks, false)); +} + +TEST(GinAllReducePolicy, YieldToDdaBySizeGfx1250) { + EXPECT_TRUE(ginAllReduceYieldToDdaBySize(countForBytes(63 * kMiB), kFloat, false, kGinAllReduceGinScatterMinBytes)); + EXPECT_FALSE(ginAllReduceYieldToDdaBySize(countForBytes(64 * kMiB), kFloat, false, kGinAllReduceGinScatterMinBytes)); +} + +TEST(GinAllReducePolicy, GinScatterDispatchGfx1250From64MiB) { + EXPECT_FALSE(ginAllReduceGinScatterDispatch(true, 64ull * kMiB - 16)); + EXPECT_TRUE(ginAllReduceGinScatterDispatch(true, 64ull * kMiB)); + EXPECT_TRUE(ginAllReduceGinScatterDispatch(true, 256ull * kMiB)); + EXPECT_FALSE(ginAllReduceGinScatterDispatch(false, 64ull * kMiB)); + EXPECT_FALSE(ginAllReduceGinScatterDispatch(false, 256ull * kMiB)); +} + +TEST(GinAllReducePolicy, GinScatterLaunchFrom64MiB) { + const size_t scratch = kGinAllReduceTwoShotScratchBytes; + const size_t chunk64 = (64ull * kMiB) / kRanks; + const size_t chunk256 = (256ull * kMiB) / kRanks; + const size_t chunk512 = (512ull * kMiB) / kRanks; + EXPECT_FALSE(ginAllReduceGinScatterLaunch(64ull * kMiB - 16, chunk64, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterLaunch(64ull * kMiB, chunk64, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterLaunch(256ull * kMiB, chunk256, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterLaunch(512ull * kMiB, chunk512, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterLaunch(2ull * 1024ull * kMiB, 256ull * kMiB, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterLaunch(kGinAllReduceGinScatterMaxBytes, 512ull * kMiB, kRanks, scratch)); +} + +TEST(GinAllReducePolicy, GinScatterLaunchWhenMessageExceedsScratch) { + const size_t chunk4g = kGinAllReduceGinScatterMaxBytes / kRanks; + EXPECT_TRUE(ginAllReduceGinScatterLaunch(kGinAllReduceGinScatterMaxBytes, chunk4g, kRanks, 512ull * kMiB)); + const size_t minScratch = ginAllReduceGinScatterStagedScratchBytes(kGinAllReduceGinScatterChunkAlign, kRanks, 1); + EXPECT_TRUE(ginAllReduceGinScatterLaunch(64ull * kMiB, (64ull * kMiB) / kRanks, kRanks, minScratch)); + EXPECT_FALSE(ginAllReduceGinScatterLaunch(64ull * kMiB, (64ull * kMiB) / kRanks, kRanks, minScratch - 1)); +} + +TEST(GinAllReducePolicy, GinScatterDefaultScratchStagesAbove512MiB) { + const size_t scratch = kGinAllReduceTwoShotScratchBytes; + EXPECT_FALSE(ginAllReduceGinScatterStaged((64ull * kMiB) / kRanks, kRanks, scratch)); + EXPECT_FALSE(ginAllReduceGinScatterStaged((256ull * kMiB) / kRanks, kRanks, scratch)); + EXPECT_FALSE(ginAllReduceGinScatterStaged((512ull * kMiB) / kRanks, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterStaged((1024ull * kMiB) / kRanks, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterStaged((2ull * 1024ull * kMiB) / kRanks, kRanks, scratch)); + EXPECT_TRUE(ginAllReduceGinScatterStaged(kGinAllReduceGinScatterMaxBytes / kRanks, kRanks, scratch)); +} + +TEST(GinAllReducePolicy, GinScatterStagesWhenMessageExceedsScratch) { + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + const size_t smallSeq = ginAllReduceGinScatterStagedScratchBytes(kGinAllReduceGinScatterChunkBytes, kRanks, 1); + ginAllReduceGinScatterChunkPlan(32ull * kMiB, kRanks, smallSeq, &nChunks, &uniformBytes, &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 2); + EXPECT_EQ(nPhases, 1); + EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); + const size_t smallPipe = ginAllReduceGinScatterStagedScratchBytes(kGinAllReduceGinScatterChunkBytes, kRanks, 2); + ginAllReduceGinScatterChunkPlan(512ull * kMiB, kRanks, smallPipe, &nChunks, &uniformBytes, &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 32); + EXPECT_EQ(nPhases, 2); + EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); +} + +TEST(GinAllReducePolicy, GinScatterNoPipelineBelow512MiB) { + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + const size_t scratch = kGinAllReduceTwoShotScratchBytes; + ginAllReduceGinScatterChunkPlan((64ull * kMiB) / kRanks, kRanks, scratch, &nChunks, &uniformBytes, &lastBytes, + &nPhases); + EXPECT_EQ(nChunks, 1); + EXPECT_EQ(nPhases, 1); + EXPECT_EQ(uniformBytes, 8ull * kMiB); + ginAllReduceGinScatterChunkPlan(32ull * kMiB, kRanks, scratch, &nChunks, &uniformBytes, &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 1); + EXPECT_EQ(nPhases, 1); + EXPECT_EQ(uniformBytes, 32ull * kMiB); + ginAllReduceGinScatterChunkPlan(63ull * kMiB, kRanks, scratch, &nChunks, &uniformBytes, &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 1); + EXPECT_EQ(nPhases, 1); +} + +TEST(GinAllReducePolicy, GinScatterChunksOverlapAt512MiB) { + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + ginAllReduceGinScatterChunkPlan(64ull * kMiB, kRanks, kGinAllReduceTwoShotScratchBytes, &nChunks, &uniformBytes, + &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 4); + EXPECT_EQ(nPhases, 1); + EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); + EXPECT_EQ(lastBytes, kGinAllReduceGinScatterChunkBytes); +} + +TEST(GinAllReducePolicy, GinScatterStagedPingPongAt1GiB) { + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + ginAllReduceGinScatterChunkPlan(128ull * kMiB, kRanks, kGinAllReduceTwoShotScratchBytes, &nChunks, &uniformBytes, + &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 4); + EXPECT_EQ(nPhases, 2); + EXPECT_EQ(uniformBytes, 32ull * kMiB); + EXPECT_EQ(lastBytes, 32ull * kMiB); +} + +TEST(GinAllReducePolicy, GinScatterSingleChunkWhenAtMostChunkSize) { + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + ginAllReduceGinScatterChunkPlan(kGinAllReduceGinScatterChunkBytes, kRanks, kGinAllReduceTwoShotScratchBytes, &nChunks, + &uniformBytes, &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 1); + EXPECT_EQ(nPhases, 1); + EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); + ginAllReduceGinScatterChunkPlan(0, kRanks, kGinAllReduceTwoShotScratchBytes, &nChunks, &uniformBytes, &lastBytes, + &nPhases); + EXPECT_EQ(nChunks, 1); + EXPECT_EQ(uniformBytes, 0u); +} + +TEST(GinAllReducePolicy, GinScatterChunksTilePerRankBytes) { + const size_t perRank = 132ull * kMiB; + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + ginAllReduceGinScatterChunkPlan(perRank, kRanks, kGinAllReduceTwoShotScratchBytes, &nChunks, &uniformBytes, &lastBytes, + &nPhases); + EXPECT_GT(nChunks, 1); + EXPECT_EQ(nChunks, 5); + EXPECT_EQ(nPhases, 2); + EXPECT_EQ(uniformBytes, 32ull * kMiB); + EXPECT_EQ(lastBytes, 4ull * kMiB); + EXPECT_EQ(uniformBytes * static_cast(nChunks - 1) + lastBytes, perRank); + EXPECT_EQ(ginAllReduceGinScatterChunkOff(nChunks - 1, uniformBytes), + uniformBytes * static_cast(nChunks - 1)); + EXPECT_EQ(ginAllReduceGinScatterChunkSize(0, nChunks, uniformBytes, lastBytes), uniformBytes); + EXPECT_EQ(ginAllReduceGinScatterChunkSize(nChunks - 1, nChunks, uniformBytes, lastBytes), lastBytes); +} + +TEST(GinAllReducePolicy, GinScatterScratchIsIncomingSlots) { + EXPECT_EQ(ginAllReduceGinScatterScratchBytes(32ull * kMiB, kRanks), 8ull * 32ull * kMiB); + EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(8ull * kMiB, kRanks, 2), 16ull * 8ull * kMiB); + EXPECT_EQ(ginAllReduceGinScatterScratchBytes(16, 1), 16u); + EXPECT_EQ(ginAllReduceGinScatterScratchBytes(16, 0), 0u); + EXPECT_EQ(ginAllReduceGinScatterFitSlot(kGinAllReduceTwoShotScratchBytes, kRanks, 1), 64ull * kMiB); + EXPECT_EQ(ginAllReduceGinScatterFitSlot(kGinAllReduceTwoShotScratchBytes, kRanks, 2), 32ull * kMiB); +} + } // namespace From 076e36ead4bba0980d46176d834377e2cee5b352 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Fri, 9 Oct 2026 09:31:26 -0500 Subject: [PATCH 2/7] address reviews --- projects/rccl/CHANGELOG.md | 2 + projects/rccl/cmake/DeviceLinker.cmake | 2 +- .../rccl/docs/api-reference/env-variables.rst | 11 + projects/rccl/src/CMakeLists.txt | 1 + .../rccl/src/algorithms/gin/gin_all_reduce.h | 12 +- .../algorithms/gin/gin_all_reduce_policy.h | 27 ++- .../algorithms/gin/sdma/all_reduce_gin_sdma.h | 201 ++++++------------ .../gin/sdma/gin_all_reduce_sdma.cu | 93 +++++--- .../algorithms/gin/sdma/gin_scatter_reduce.h | 112 ++++++++++ projects/rccl/src/gin/README.md | 4 +- projects/rccl/src/rccl_wrap.cc | 5 +- projects/rccl/test/CMakeLists.txt | 26 +++ .../test/gin/gin_all_reduce_policy_test.cpp | 28 ++- .../test/gin/gin_scatter_reduce_gpu_test.cpp | 128 +++++++++++ 14 files changed, 457 insertions(+), 195 deletions(-) create mode 100644 projects/rccl/src/algorithms/gin/sdma/gin_scatter_reduce.h create mode 100644 projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 10f5cd492240..b1826eef8045 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -22,6 +22,8 @@ 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. +* `NCCL_GIN_ALLREDUCE_GFX1250_ENABLE` (default `0`): opt in to GinScatter AllReduce on gfx1250 for symmetric sum messages from 64 MiB through 4 GiB. Left off, gfx1250 keeps the pre-existing AllReduce selectors. +* `NCCL_GIN_ALLREDUCE_SCRATCH_BYTES` (default 512 MiB): GinScatter incoming scratch. `0` skips the allocation and GinScatter is not selected. The buffer is allocated only on the GinScatter path. ### 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`. diff --git a/projects/rccl/cmake/DeviceLinker.cmake b/projects/rccl/cmake/DeviceLinker.cmake index 78acd876228d..0d46146e9aaf 100644 --- a/projects/rccl/cmake/DeviceLinker.cmake +++ b/projects/rccl/cmake/DeviceLinker.cmake @@ -923,7 +923,7 @@ endif() # =========================================================================== set(GIN_ROCSHMEM_QP_BC_FLAG "") set(GIN_ROCSHMEM_QP_DEPS "") -if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR AND DL_GPU_TARGETS) +if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR AND DL_GPU_TARGETS AND GENERATE_SYM_KERNELS) list(GET DL_GPU_TARGETS 0 _bc_arch) set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc") find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED) diff --git a/projects/rccl/docs/api-reference/env-variables.rst b/projects/rccl/docs/api-reference/env-variables.rst index 1b730f2e265f..7bbde0767d4e 100644 --- a/projects/rccl/docs/api-reference/env-variables.rst +++ b/projects/rccl/docs/api-reference/env-variables.rst @@ -511,6 +511,17 @@ in the following table. | ``6``: rocSHMEM GDA (required for that backend to initialize) | ``7``: Anvil SDMA (also accepts unset) + * - | ``NCCL_GIN_ALLREDUCE_GFX1250_ENABLE`` + | Opt in to GinScatter AllReduce on gfx1250. + - | ``0``: gfx1250 keeps the pre-existing AllReduce selectors (default) + | ``1``: symmetric ``ncclSum`` fp32/fp16/bf16 messages from 64 MiB + through 4 GiB use GinScatter when scratch can hold them + + * - | ``NCCL_GIN_ALLREDUCE_SCRATCH_BYTES`` + | GinScatter incoming scratch size, in bytes. + - | ``536870912`` (512 MiB) by default, matching the CE AllReduce pool + | ``0``: do not allocate scratch, and do not select GinScatter + * - | ``NCCL_RMA_PLUGIN`` | Selects external one-sided RMA plugins, which are also the backend the built-in GIN proxy forwards to. diff --git a/projects/rccl/src/CMakeLists.txt b/projects/rccl/src/CMakeLists.txt index cb250b9972e2..663b947bf21c 100644 --- a/projects/rccl/src/CMakeLists.txt +++ b/projects/rccl/src/CMakeLists.txt @@ -164,6 +164,7 @@ set(SRC_FILES algorithms/gin/gin_all_reduce.h algorithms/gin/gin_all_reduce_policy.h algorithms/gin/sdma/all_reduce_gin_sdma.h + algorithms/gin/sdma/gin_scatter_reduce.h include/alloc.h include/allocator.h include/alltoallv_meta.h diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce.h b/projects/rccl/src/algorithms/gin/gin_all_reduce.h index f261cdbf35d6..28373f4a54d2 100755 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce.h @@ -3,8 +3,8 @@ * * Host entry points for the GIN-SDMA AllReduce path selected by rcclSelectAllReduce * when symmetric windows are used. By default gfx950 takes GIN two-shot for - * messages >= 256 MiB; gfx1250 with Anvil SDMA takes GinScatter from 64 MiB. - * Smaller messages use DDA AllReduce. + * messages >= 256 MiB. gfx1250 GinScatter (64 MiB to 4 GiB) is opt-in via + * NCCL_GIN_ALLREDUCE_GFX1250_ENABLE. Smaller messages use DDA AllReduce. * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables LSA one-shot (<= 4 MiB) and * LSA two-shot ((4 MiB, GIN floor)). * See LICENSE.txt for license information. @@ -27,7 +27,9 @@ struct ncclDevrWindow; // LSA one-shot for messages <= kGinAllReduceLsaOneShotMaxBytes (force-enable only). // LSA two-shot between 4 MiB and the GIN floor (force-enable only). // gfx950 + Anvil SDMA: GIN two-shot for messages >= kGinAllReduceGinTwoShotMinBytes (256 MiB). -// gfx1250 + Anvil SDMA: GinScatter for messages >= kGinAllReduceGinScatterMinBytes (64 MiB). +// gfx1250 + Anvil SDMA: GinScatter for messages in +// [kGinAllReduceGinScatterMinBytes, kGinAllReduceGinScatterMaxBytes] when +// NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1. // Lazily created on the first eligible AllReduce and torn down with the comm. // Declared unconditionally: ncclComm embeds this even when ENABLE_ROCSHMEM_GIN is off. @@ -40,8 +42,8 @@ struct ncclDevrWindow; // intraGpuCtaBar is two device uint32s (arrived, sense) for the GIN two-shot intra-GPU // CTA barrier; the pointer is stable, the words themselves advance on the device. // scratch / scratchWin are the symmetric staging buffer for GinScatter (512 MiB of -// CE-style ping-pong incoming slots). Allocated and window-registered once on -// gfx1250 before GIN connect. +// CE-style ping-pong incoming slots by default). Allocated on the GinScatter path +// only, before GIN connect when that path is the first GIN AllReduce. struct ncclGinAllReduceState { bool initialized; struct ncclDevComm devComm; diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h index bc576ce6a3a3..cd0f084f8027 100644 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h @@ -7,7 +7,7 @@ * host+device so the GinScatter kernel splits the same way the tests do. * * gfx950 default: GIN two-shot from 256 MiB (LSA RS + GIN AG, Anvil SDMA). - * gfx1250 + Anvil SDMA: GinScatter from 64 MiB. + * gfx1250 + Anvil SDMA: GinScatter from 64 MiB when NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1. * GinScatter scratch follows CE AllReduce (ce_reduce_impl.h / ce_coll.h): * 512 MiB of incoming ping-pong slots, reduced output in the user recvbuff. * See LICENSE.txt for license information. @@ -61,11 +61,16 @@ GIN_AR_HD size_t ginAllReduceGinScatterAlignDown(size_t n, size_t align) { } // CE ceARTmpBuf: nSlots * nRanks * slotBytes. Reduced output is the user recvbuff. +// The two-argument form is one incoming column per rank. Ping-pong passes nSlots=2. GIN_AR_HD size_t ginAllReduceGinScatterStagedScratchBytes(size_t slotBytes, int nRanks, int nSlots) { if (nRanks <= 0 || nSlots <= 0) return 0; return static_cast(nSlots) * static_cast(nRanks) * slotBytes; } +GIN_AR_HD size_t ginAllReduceGinScatterStagedScratchBytes(size_t slotBytes, int nRanks) { + return ginAllReduceGinScatterStagedScratchBytes(slotBytes, nRanks, 1); +} + // Largest 16-byte-aligned per-rank slot that fits in scratch with nSlots windows. GIN_AR_HD size_t ginAllReduceGinScatterFitSlot(size_t scratchBytes, int nRanks, int nSlots) { if (nRanks <= 0 || nSlots <= 0) return 0; @@ -79,7 +84,7 @@ GIN_AR_HD bool ginAllReduceGinScatterStaged(size_t perRankBytes, int nRanks, siz return scratchBytes < static_cast(nRanks) * perRankBytes; } -// Split one per-rank slice. Full-slot path: one chunk below 512 MiB, 16 MiB +// Split one per-rank slice. Full-slot path: one chunk below 512 MiB, 64 MiB // stages at 512 MiB+. Staged path (message > scratch): CE ping-pong — two slots // and overlap at >= 512 MiB; one slot and sequential reuse below 512 MiB. GIN_AR_HD void ginAllReduceGinScatterChunkPlan(size_t perRankBytes, int nRanks, size_t scratchBytes, int* nChunks, @@ -148,13 +153,6 @@ GIN_AR_HD size_t ginAllReduceGinScatterChunkSize(int chunk, int nChunks, size_t #undef GIN_AR_HD -// Full-slot incoming columns (no reduced scratch column; CE writes the reduced -// shard to recvbuff). Used to decide whether the message fits without staging. -inline size_t ginAllReduceGinScatterScratchBytes(size_t chunkBytes, int nRanks) { - if (nRanks <= 0) return 0; - return static_cast(nRanks) * chunkBytes; -} - // Two-shot kernels require a whole number of elements per rank and a 16-byte // per-rank slice so vector loads stay aligned. inline bool ginAllReduceTwoShotEligible(size_t count, size_t typeSize, int nRanks) { @@ -173,9 +171,9 @@ inline bool ginAllReduceGinTwoShotEligible(size_t count, size_t typeSize, int nR // ginScatter from 64 MiB when scratch holds either nRanks full incoming columns // or at least one 16-byte-aligned staging slot per rank. inline bool ginAllReduceGinScatterLaunch(size_t bytes, size_t chunkBytes, int nRanks, size_t scratchBytes) { - if (bytes < kGinAllReduceGinScatterMinBytes) return false; + if (bytes < kGinAllReduceGinScatterMinBytes || bytes > kGinAllReduceGinScatterMaxBytes) return false; if (nRanks <= 0) return false; - if (scratchBytes >= ginAllReduceGinScatterScratchBytes(chunkBytes, nRanks)) return true; + if (scratchBytes >= ginAllReduceGinScatterStagedScratchBytes(chunkBytes, nRanks)) return true; return ginAllReduceGinScatterFitSlot(scratchBytes, nRanks, 1) >= kGinAllReduceGinScatterChunkAlign; } @@ -204,10 +202,11 @@ inline bool ginAllReduceYieldToDdaBySize(size_t count, size_t typeSize, bool for return count * typeSize < ginMinBytes; } -// Host dispatch for ginAllReduceTwoShotGinScatterKernel: gfx1250 and >= 64 MiB. -// Eligibility still requires the Anvil SDMA GIN backend before launch. +// Host dispatch for ginAllReduceTwoShotGinScatterKernel: gfx1250 and +// [64 MiB, 4 GiB]. Eligibility still requires the Anvil SDMA GIN backend, +// NCCL_GIN_ALLREDUCE_GFX1250_ENABLE, and a scratch pool that can hold the message. inline bool ginAllReduceGinScatterDispatch(bool isGfx1250, size_t bytes) { - return isGfx1250 && bytes >= kGinAllReduceGinScatterMinBytes; + return isGfx1250 && bytes >= kGinAllReduceGinScatterMinBytes && bytes <= kGinAllReduceGinScatterMaxBytes; } #endif diff --git a/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h b/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h index 2be0cad01342..69bec2a3b2de 100755 --- a/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h +++ b/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h @@ -9,7 +9,8 @@ * (self from send, remotes from scratch). Scratch is 512 MiB of CE-style * ping-pong incoming slots; reduced output is recvbuff. Local reduce matches * CE (16B vectors, GpuUnroll=4, ranks 0..nRanks-1). Used on gfx1250 from 64 MiB - * when the Anvil SDMA GIN backend is enabled. + * to 4 GiB when NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1 and the Anvil SDMA GIN + * backend is enabled. * * One-shot follows projects/rccl-tests/src/all_reduce.cu allReduceLsaKernel. ******************************************************************************/ @@ -20,6 +21,7 @@ #include "nccl_device/ptr.h" #include "algorithms/dda/device/CollCommon.h" #include "algorithms/gin/gin_all_reduce_policy.h" +#include "algorithms/gin/sdma/gin_scatter_reduce.h" namespace gin::sdma { using dda::common::vecElementAdd; @@ -67,33 +69,6 @@ __launch_bounds__(512) bar.sync(ncclCoopCta(), cuda::memory_order_release); } -// 16B window accesses through the global aperture. -// -// ncclGetLsaPointer hands back a generic (flat) pointer built by bit-twiddling a base loaded from -// the window struct, so nothing downstream can prove it is device memory and the vector accesses -// compile to flat_load_dwordx4 / flat_store_dwordx4. Peer windows live in the global aperture, so -// casting to an address_space(1) pointer leaves the semantics untouched (plain, non-atomic, ordered -// by the surrounding LSA barriers) while emitting global_load_dwordx4 / global_store_dwordx4 — -// exactly what rccl_ptr.h prescribes for hot paths. If an ISA dump of this kernel mentions "flat", -// one of these casts was lost. -__device__ __forceinline__ uint4 lsaLoadVec(const char* p) { - union { - v4u vec; - uint4 val; - } u; - u.vec = *(v4u_gptr)(const_cast(p)); - return u.val; -} - -__device__ __forceinline__ void lsaStoreVec(char* p, uint4 val) { - union { - v4u vec; - uint4 val; - } u; - u.val = val; - *(v4u_gptr)(p) = u.vec; -} - // A rank's own column of an LSA window, resolved for every peer at once. // // ncclGetLsaPointer(win, off, peer) is lsaFlatBase + (peer * stride4G << 32) + off, i.e. linear in @@ -336,11 +311,23 @@ __launch_bounds__(512) } +// Anvil SDMA WeakSignalInc adds into one word per index. A shared index is not a +// barrier: one peer's later chunk can reach an aggregate nRanks-1 count before a +// slower peer's chunk arrives. Scatter and slot-free each use one index per source +// rank (same layout as ncclGinBarrierSession). All-gather issues one put per peer, +// so a single counter is enough. Indices stay inside the ginSignalCount reserved +// at init (kGinAllReduceLsaCtas). +constexpr unsigned kGinScatterSignalBase = 0; +constexpr unsigned kGinScatterDoneSignalBase = static_cast(kGinAllReduceMaxRanks); +constexpr unsigned kGinScatterAgSignal = 2u * static_cast(kGinAllReduceMaxRanks); + // GIN-put one pipeline chunk of send[dst] into dest dst's staging slot for every dst != rank. +// The destination waits on this rank's own scatter signal. __device__ __forceinline__ void ginScatterPutRemoteChunk(ncclGin& gin, ncclTeam world, int rank, int nRanks, ncclWindow_t scratchWin, size_t destSlotOff, ncclWindow_t sendWin, size_t sendOff, size_t perRankBytes, - size_t sendChunkOff, size_t chunkBytes, unsigned signalIndex) { + size_t sendChunkOff, size_t chunkBytes) { + const unsigned signalIndex = kGinScatterSignalBase + static_cast(rank); for (int dst = static_cast(threadIdx.x); dst < nRanks; dst += static_cast(blockDim.x)) { if (dst == rank) { continue; @@ -368,88 +355,18 @@ __device__ __forceinline__ const char* ginScatterIncomingBase(const char* scratc return scratchBase + static_cast(slot) * static_cast(nRanks) * slotBytes; } -// CE ncclCeLocalReduceKernelVec: 16B vectors, GpuUnroll=4, rank loop 0..nRanks-1 -// into recvbuff + rank*shard. Self is not written to scratch (GIN-put skips it), -// so rank == myRank loads from send instead of the incoming slot. -template -__device__ __forceinline__ const char* ginScatterReduceSrc(int peer, int rank, const char* selfSend, - const char* incomingBase, size_t incomingStride, size_t elem, - size_t elemStart, bool staged) { - if (peer == rank) { - return selfSend + elem * sizeof(T); - } - const size_t srcOff = staged ? (elem - elemStart) * sizeof(T) : elem * sizeof(T); - return incomingBase + static_cast(peer) * incomingStride + srcOff; -} - -template -__device__ __forceinline__ void ginScatterReduceChunk(const char* selfSend, const char* incomingBase, char* reducedOut, - int rank, int nRanks, size_t incomingStride, size_t elemStart, - size_t elemEnd, bool staged) { - constexpr int W = static_cast(sizeof(uint4) / sizeof(T)); - constexpr int U = 4; - const size_t nVec = (elemEnd - elemStart) / static_cast(W); - const size_t tid = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; - const size_t stride = static_cast(blockDim.x) * gridDim.x; - - size_t vi = tid; - for (; vi + static_cast(U - 1) * stride < nVec; vi += stride * static_cast(U)) { - uint4 acc[U]; - size_t elem[U]; -#pragma unroll - for (int i = 0; i < U; ++i) { - elem[i] = elemStart + (vi + static_cast(i) * stride) * static_cast(W); - acc[i] = lsaLoadVec(ginScatterReduceSrc(0, rank, selfSend, incomingBase, incomingStride, elem[i], elemStart, - staged)); - } -#pragma clang loop unroll(disable) vectorize(disable) - for (int r = 1; r < nRanks; ++r) { - uint4 tmp[U]; -#pragma unroll - for (int i = 0; i < U; ++i) { - tmp[i] = lsaLoadVec(ginScatterReduceSrc(r, rank, selfSend, incomingBase, incomingStride, elem[i], elemStart, - staged)); - } -#pragma unroll - for (int i = 0; i < U; ++i) { - acc[i] = vecElementAdd(acc[i], tmp[i]); - } - } -#pragma unroll - for (int i = 0; i < U; ++i) { - lsaStoreVec(reducedOut + elem[i] * sizeof(T), acc[i]); - } - } - - const size_t vectorTailStart = nVec - (nVec % (stride * static_cast(U))); - if (nVec > 0 && nVec % (stride * static_cast(U)) != 0) { - for (size_t vIdx = vectorTailStart + tid; vIdx < nVec; vIdx += stride) { - const size_t e = elemStart + vIdx * static_cast(W); - uint4 acc = - lsaLoadVec(ginScatterReduceSrc(0, rank, selfSend, incomingBase, incomingStride, e, elemStart, staged)); -#pragma clang loop unroll(disable) vectorize(disable) - for (int r = 1; r < nRanks; ++r) { - acc = vecElementAdd( - acc, lsaLoadVec(ginScatterReduceSrc(r, rank, selfSend, incomingBase, incomingStride, e, elemStart, - staged))); - } - lsaStoreVec(reducedOut + e * sizeof(T), acc); - } - } -} - __device__ __forceinline__ void ginScatterIssuePut(ncclGin& gin, ncclTeam world, int rank, int nRanks, ncclWindow_t scratchWin, size_t scratchOff, ncclWindow_t sendWin, size_t sendOff, size_t perRankBytes, size_t slotBytes, size_t uniformBytes, size_t lastBytes, int nChunks, int nPhases, - bool staged, int chunk, unsigned signalIndex) { + bool staged, int chunk) { const size_t chunkOff = ginAllReduceGinScatterChunkOff(chunk, uniformBytes); const size_t chunkBytes = ginAllReduceGinScatterChunkSize(chunk, nChunks, uniformBytes, lastBytes); const int slot = staged ? ((nPhases <= 1) ? 0 : (chunk % nPhases)) : 0; const size_t destOff = ginScatterPutDestOff(scratchOff, rank, nRanks, perRankBytes, slotBytes, slot, staged ? 0 : chunkOff, staged); ginScatterPutRemoteChunk(gin, world, rank, nRanks, scratchWin, destOff, sendWin, sendOff, perRankBytes, chunkOff, - chunkBytes, signalIndex); + chunkBytes); } // Thread 0 polls; callers that need the rest of the CTA to observe completion @@ -460,10 +377,22 @@ __device__ __forceinline__ void ginScatterWaitThread0(ncclGin& gin, unsigned sig } } +// Thread 0 waits until every other rank's own signal has advanced by `add`. +// `baseline` is this thread's snapshot and is only valid on thread 0. +__device__ __forceinline__ void ginScatterWaitPeersThread0(ncclGin& gin, unsigned base, const uint64_t* baseline, + int rank, int nRanks, uint64_t add) { + if (threadIdx.x != 0) return; + for (int peer = 0; peer < nRanks; ++peer) { + if (peer == rank) continue; + gin.waitSignal(ncclCoopThread(), base + static_cast(peer), baseline[peer] + add); + } +} + // One-way doorbell: this rank finished reading a reused scratch slot. Cheaper than -// a world barrier (CE uses per-slot signal==0 the same way). -__device__ __forceinline__ void ginScatterNotifySlotFree(ncclGin& gin, ncclTeam world, int rank, int nRanks, - unsigned doneSignal) { +// a world barrier (CE uses per-slot signal==0 the same way). Each source uses its +// own done-signal index so a fast peer cannot release a slot a slow peer still reads. +__device__ __forceinline__ void ginScatterNotifySlotFree(ncclGin& gin, ncclTeam world, int rank, int nRanks) { + const unsigned doneSignal = kGinScatterDoneSignalBase + static_cast(rank); for (int dst = static_cast(threadIdx.x); dst < nRanks; dst += static_cast(blockDim.x)) { if (dst == rank) { continue; @@ -484,9 +413,10 @@ __device__ __forceinline__ void ginScatterNotifySlotFree(ncclGin& gin, ncclTeam // (self from send). Below 512 MiB a step completes scatter before reduce. At // 512 MiB+ put of chunk i+1 overlaps reduce of chunk i. CTA 0 GIN all-gathers. // -// Sync: one world acquire so all ranks snapshot GIN signal baselines. Each pipeline -// step is waitSignal on fused WeakSignalInc (scatter/AG) or a one-way slot-free -// doorbell — not ginBar.sync or gin.flush. Intra-GPU CTA bars only join the 56 CTAs. +// Sync: one world acquire so all ranks snapshot per-peer GIN signal baselines. +// Each pipeline step waits on that peer's scatter or slot-free signal. All-gather +// is one put per peer on a shared counter, then gin.flush so recvbuff can be reused. +// Intra-GPU CTA bars only join the 56 CTAs. template #if defined(USE_ROCM) __launch_bounds__(512) @@ -496,12 +426,9 @@ __launch_bounds__(512) size_t scratchOff, size_t countPerRank, int nRanks, size_t scratchBytes, uint32_t* intraGpuCtaBar) { constexpr int ginContext = 0; - constexpr unsigned scatterSignal = 0; - constexpr unsigned agSignal = 1; - constexpr unsigned doneSignal = 2; // gin/ginBar must outlive the first put; only CTA 0 snapshots signals and world-syncs. - // Pipeline steps use waitSignal / WeakSignalInc (CE doorbells), not world barriers. + // Pipeline steps use per-peer waitSignal / WeakSignalInc, not a shared counter. ncclGin gin{devComm, ginContext}; ncclTeam world = ncclTeamWorld(devComm); ncclCoopCta cta; @@ -524,36 +451,41 @@ __launch_bounds__(512) const size_t incomingStride = staged ? slotBytes : perRankBytes; const bool overlap = (nChunks > 1) && !(staged && nPhases <= 1); - const uint64_t signalsPerChunk = static_cast(nRanks - 1); - uint64_t scatterValue = 0; - uint64_t doneValue = 0; + uint64_t scatterBase[kGinAllReduceMaxRanks] = {}; + uint64_t doneBase[kGinAllReduceMaxRanks] = {}; uint64_t agValue = 0; - // One world acquire so every rank snapshots signal baselines before any put. - // Later pipeline steps are waitSignal doorbells (CE), not ginBar.sync. + // One world acquire so every rank snapshots per-peer baselines before any put. + // Later pipeline steps are per-peer waitSignal doorbells, not ginBar.sync. if (blockIdx.x == 0) { - scatterValue = gin.readSignal(scatterSignal); - doneValue = gin.readSignal(doneSignal); - agValue = gin.readSignal(agSignal); + if (threadIdx.x == 0) { + const int peers = nRanks < kGinAllReduceMaxRanks ? nRanks : kGinAllReduceMaxRanks; + for (int peer = 0; peer < peers; ++peer) { + scatterBase[peer] = gin.readSignal(kGinScatterSignalBase + static_cast(peer)); + doneBase[peer] = gin.readSignal(kGinScatterDoneSignalBase + static_cast(peer)); + } + agValue = gin.readSignal(kGinScatterAgSignal); + } ginBar.sync(cta, cuda::memory_order_acquire, ncclGinFenceLevel::None); } if (overlap) { if (blockIdx.x == 0) { ginScatterIssuePut(gin, world, devComm.rank, nRanks, scratchWin, scratchOff, sendWin, sendOff, perRankBytes, - slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, 0, scatterSignal); - ginScatterWaitThread0(gin, scatterSignal, scatterValue + signalsPerChunk); + slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, 0); + ginScatterWaitPeersThread0(gin, kGinScatterSignalBase, scatterBase, devComm.rank, nRanks, 1); } ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); for (int c = 0; c < nChunks; ++c) { if (blockIdx.x == 0 && c + 1 < nChunks) { if (staged && c + 1 >= nPhases) { - ginScatterWaitThread0(gin, doneSignal, doneValue + static_cast(c + 2 - nPhases) * signalsPerChunk); + ginScatterWaitPeersThread0(gin, kGinScatterDoneSignalBase, doneBase, devComm.rank, nRanks, + static_cast(c + 2 - nPhases)); __syncthreads(); } ginScatterIssuePut(gin, world, devComm.rank, nRanks, scratchWin, scratchOff, sendWin, sendOff, perRankBytes, - slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, c + 1, scatterSignal); + slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, c + 1); } const size_t chunkOff = ginAllReduceGinScatterChunkOff(c, uniformBytes); const size_t chunkBytes = ginAllReduceGinScatterChunkSize(c, nChunks, uniformBytes, lastBytes); @@ -563,11 +495,12 @@ __launch_bounds__(512) ginScatterReduceChunk(selfSend, ginScatterIncomingBase(scratchBase, nRanks, slotBytes, slot, staged), reducedOut, devComm.rank, nRanks, incomingStride, elemStart, elemEnd, staged); if (blockIdx.x == 0 && c + 1 < nChunks) { - ginScatterWaitThread0(gin, scatterSignal, scatterValue + static_cast(c + 2) * signalsPerChunk); + ginScatterWaitPeersThread0(gin, kGinScatterSignalBase, scatterBase, devComm.rank, nRanks, + static_cast(c + 2)); } ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); if (blockIdx.x == 0 && staged && c + nPhases < nChunks) { - ginScatterNotifySlotFree(gin, world, devComm.rank, nRanks, doneSignal); + ginScatterNotifySlotFree(gin, world, devComm.rank, nRanks); __syncthreads(); } } @@ -575,12 +508,14 @@ __launch_bounds__(512) for (int c = 0; c < nChunks; ++c) { if (blockIdx.x == 0) { if (c > 0 && staged) { - ginScatterWaitThread0(gin, doneSignal, doneValue + static_cast(c) * signalsPerChunk); + ginScatterWaitPeersThread0(gin, kGinScatterDoneSignalBase, doneBase, devComm.rank, nRanks, + static_cast(c)); __syncthreads(); } ginScatterIssuePut(gin, world, devComm.rank, nRanks, scratchWin, scratchOff, sendWin, sendOff, perRankBytes, - slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, c, scatterSignal); - ginScatterWaitThread0(gin, scatterSignal, scatterValue + static_cast(c + 1) * signalsPerChunk); + slotBytes, uniformBytes, lastBytes, nChunks, nPhases, staged, c); + ginScatterWaitPeersThread0(gin, kGinScatterSignalBase, scatterBase, devComm.rank, nRanks, + static_cast(c + 1)); } ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); const size_t chunkOff = ginAllReduceGinScatterChunkOff(c, uniformBytes); @@ -592,23 +527,25 @@ __launch_bounds__(512) reducedOut, devComm.rank, nRanks, incomingStride, elemStart, elemEnd, staged); ginIntraGpuCtaBarrier(intraGpuCtaBar, static_cast(gridDim.x)); if (blockIdx.x == 0 && staged && c + 1 < nChunks) { - ginScatterNotifySlotFree(gin, world, devComm.rank, nRanks, doneSignal); + ginScatterNotifySlotFree(gin, world, devComm.rank, nRanks); __syncthreads(); } } } if (blockIdx.x == 0) { - // Shot 2: waitSignal on fused AG puts is the completion doorbell; no world barrier. + // Shot 2: one put per peer, so a shared counter matches the sibling kernels. + // flush releases recvbuff for reuse; the wait only counts inbound puts. for (int dst = static_cast(threadIdx.x); dst < nRanks; dst += static_cast(blockDim.x)) { if (dst == devComm.rank) { continue; } gin.put(world, dst, recvWin, sliceRecvByteOff, recvWin, sliceRecvByteOff, perRankBytes, - ncclGin_WeakSignalInc{agSignal}); + ncclGin_WeakSignalInc{kGinScatterAgSignal}); } - ginScatterWaitThread0(gin, agSignal, agValue + static_cast(nRanks - 1)); + ginScatterWaitThread0(gin, kGinScatterAgSignal, agValue + static_cast(nRanks - 1)); __syncthreads(); + gin.flush(cta); } } diff --git a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu index 2b084921c839..f786aabb75d8 100755 --- a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu +++ b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu @@ -3,12 +3,13 @@ * * GIN-SDMA AllReduce for single-node (scaleup-only) symmetric windows. * Default: gfx950 messages >= 256 MiB take GIN two-shot (LSA reduce-scatter + - * GIN all-gather). gfx1250 messages >= 64 MiB take GinScatter over Anvil SDMA. + * GIN all-gather). gfx1250 GinScatter is opt-in + * (NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1) for messages in [64 MiB, 4 GiB]. * Smaller messages fall through to DDA AllReduce. * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables the LSA bands: * <= 4 MiB — LSA one-shot * (4 MiB, GIN floor) — LSA two-shot - * >= GIN floor — GIN two-shot (gfx950) or GinScatter (gfx1250) + * >= GIN floor — GIN two-shot (gfx950) or GinScatter (gfx1250, when enabled) * * Compiled with NCCL_GIN_ANVIL_SDMA_ENABLE=1 and NCCL_GIN_PROXY_ENABLE=0 so * ncclGinCallImpl resolves the SDMA backend at compile time. @@ -35,8 +36,13 @@ NCCL_PARAM(GinAllReduceEnable, "GIN_ALLREDUCE_ENABLE", 1); // (256 MiB two-shot on gfx950, 64 MiB GinScatter on gfx1250). // Set to 1 to also take LSA one-shot / LSA two-shot for smaller sizes. RCCL_PARAM(GinAllReduceForceEnable, "GIN_ALLREDUCE_FORCE_ENABLE", 0); -// Symmetric 512 MiB scratch for GinScatter CE ping-pong incoming slots. -// Allocated once on gfx1250 before GIN connect. 0 skips allocation. +// gfx1250 was not GIN-eligible before GinScatter. Keep that path opt-in until it +// has soaked; 0 leaves gfx1250 on the pre-existing AllReduce selectors. +NCCL_PARAM(GinAllReduceGfx1250Enable, "GIN_ALLREDUCE_GFX1250_ENABLE", 0); +// Symmetric scratch for GinScatter CE ping-pong incoming slots. Default matches +// CE's 512 MiB pool and is divided per rank; it is not grown with the message. +// Allocated on the GinScatter path only. 0 skips allocation and GinScatter +// declines so another algorithm can run. NCCL_PARAM(GinAllReduceScratchBytes, "GIN_ALLREDUCE_SCRATCH_BYTES", static_cast(kGinAllReduceTwoShotScratchBytes)); // LSA two-shot tuning. CTAs default to kGinAllReduceLsaTwoShotCtasPerPeer * nRanks; the DDA IPC @@ -62,16 +68,17 @@ static bool ginAllReduceScratchWinReady(ncclComm* comm) { static ncclResult_t ginAllReduceFreeScratch(ncclComm* comm) { struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + ncclResult_t ret = ncclSuccess; if (state->scratchWin != nullptr) { - NCCLCHECK(ncclCommWindowDeregister(comm, state->scratchWin->vidmem)); + NCCLCHECKIGNORE(ncclCommWindowDeregister(comm, state->scratchWin->vidmem), ret); state->scratchWin = nullptr; } if (state->scratch != nullptr) { - NCCLCHECK(ncclMemFree(state->scratch)); + NCCLCHECKIGNORE(ncclMemFree(state->scratch), ret); state->scratch = nullptr; } state->scratchBytes = 0; - return ncclSuccess; + return ret; } static ncclResult_t ginAllReduceScratchInitOnce(ncclComm* comm) { @@ -88,16 +95,25 @@ static ncclResult_t ginAllReduceScratchInitOnce(ncclComm* comm) { void* scratch = nullptr; ncclWindow_t scratchWinDev = nullptr; - NCCLCHECK(ncclMemAlloc(&scratch, allocBytes)); - NCCLCHECK(ncclDevrWindowRegisterInGroup(comm, scratch, allocBytes, NCCL_WIN_COLL_SYMMETRIC, &scratchWinDev)); - NCCLCHECK(ncclDevrFindWindow(comm, scratch, &state->scratchWin)); - if (scratchWinDev == nullptr || !ginAllReduceScratchWinReady(comm)) { + ncclResult_t ret = ncclMemAlloc(&scratch, allocBytes); + if (ret != ncclSuccess) return ret; + + ret = ncclDevrWindowRegisterInGroup(comm, scratch, allocBytes, NCCL_WIN_COLL_SYMMETRIC, &scratchWinDev); + if (ret != ncclSuccess) { + (void)ncclMemFree(scratch); + return ret; + } + + ret = ncclDevrFindWindow(comm, scratch, &state->scratchWin); + if (ret != ncclSuccess || scratchWinDev == nullptr || !ginAllReduceScratchWinReady(comm)) { if (scratchWinDev != nullptr) { (void)ncclCommWindowDeregister(comm, scratchWinDev); + } else if (state->scratchWin != nullptr && state->scratchWin->vidmem != nullptr) { + (void)ncclCommWindowDeregister(comm, state->scratchWin->vidmem); } state->scratchWin = nullptr; - NCCLCHECK(ncclMemFree(scratch)); - return ncclInternalError; + (void)ncclMemFree(scratch); + return ret != ncclSuccess ? ret : ncclInternalError; } state->scratch = scratch; state->scratchBytes = allocBytes; @@ -121,11 +137,6 @@ static ncclResult_t ginAllReduceScratchInitOnce(ncclComm* comm) { // barrier before every rank has finished initializing. static ncclResult_t ncclGinAllReduceInitOnce(ncclComm* comm) { NCCLCHECK(ncclDevrInitOnce(comm)); - // Window-register GinScatter scratch while GIN is still off so - // ncclDevrCommCreateInternal GIN-registers every memHead entry, including this one. - if (ginAllReduceIsGfx1250(comm)) { - NCCLCHECK(ginAllReduceScratchInitOnce(comm)); - } struct ncclGinAllReduceState* state = &comm->ginAllReduceState; if (state->initialized) { return ncclSuccess; @@ -134,6 +145,7 @@ static ncclResult_t ncclGinAllReduceInitOnce(ncclComm* comm) { struct ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER; reqs.lsaBarrierCount = kGinAllReduceLsaTwoShotMaxCtas; // ginAllReduceTwoShotKernel: one world barrier + GIN signal per CTA. + // GinScatter needs 2 * kGinAllReduceMaxRanks + 1 per-peer signals, which fits here. reqs.barrierCount = kGinAllReduceLsaCtas; reqs.ginSignalCount = kGinAllReduceLsaCtas; reqs.ginConnectionType = NCCL_GIN_CONNECTION_FULL; @@ -274,6 +286,10 @@ static ncclResult_t ncclAllReduceGinSdmaGinScatterTyped(const void* sendbuff, vo ncclComm* comm, cudaStream_t stream, struct ncclDevrWindow* sendWin, struct ncclDevrWindow* recvWin) { + // Register scratch before the first GIN connect when this is the first GIN + // AllReduce. A later GinScatter still registers: window registration + // GIN-assigns memories once ginEnabled is set. LSA and two-shot do not pay it. + NCCLCHECK(ginAllReduceScratchInitOnce(comm)); NCCLCHECK(ncclGinAllReduceInitOnce(comm)); const size_t sendOff = @@ -294,10 +310,11 @@ static ncclResult_t ncclAllReduceGinSdmaGinScatterTyped(const void* sendbuff, vo return ncclInvalidUsage; } - gin::sdma::ginAllReduceTwoShotGinScatterKernel - <<>>( - state->devComm, sendWin->vidmem, sendOff, recvWin->vidmem, recvOff, state->scratchWin->vidmem, - /*scratchOff=*/0, countPerRank, comm->nRanks, state->scratchBytes, state->intraGpuCtaBar); + const hipEvent_t stopEvent = rcclTakeAddonStopEvent(comm); + hipExtLaunchKernelGGL((gin::sdma::ginAllReduceTwoShotGinScatterKernel), kGinAllReduceLsaCtas, + kGinAllReduceLsaThreadsPerCta, 0, stream, /*startEvent=*/nullptr, stopEvent, /*flags=*/0, + state->devComm, sendWin->vidmem, sendOff, recvWin->vidmem, recvOff, state->scratchWin->vidmem, + /*scratchOff=*/0, countPerRank, comm->nRanks, state->scratchBytes, state->intraGpuCtaBar); CUDACHECK(cudaGetLastError()); return ncclSuccess; } @@ -330,7 +347,11 @@ static bool ginAllReduceBaseEligible(ncclComm* comm, const void* sendbuff, void* if (comm == nullptr || sendbuff == nullptr || recvbuff == nullptr) return false; if (count == 0) return false; - if (!IsArchMatch(comm->archName, "gfx950") && !ginAllReduceIsGfx1250(comm)) return false; + if (ginAllReduceIsGfx1250(comm)) { + if (ncclParamGinAllReduceGfx1250Enable() == 0) return false; + } else if (!IsArchMatch(comm->archName, "gfx950")) { + return false; + } if (op != ncclSum) return false; if (datatype != ncclFloat32 && datatype != ncclFloat16 && datatype != ncclBfloat16) return false; if (!comm->symmetricSupport) return false; @@ -357,8 +378,21 @@ static bool ginAllReduceBaseEligible(ncclComm* comm, const void* sendbuff, void* bool ncclAllReduceGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op) { if (!ginAllReduceBaseEligible(comm, sendbuff, recvbuff, count, datatype, op)) return false; - return ginAllReduceSizePolicyEligible(count, ncclTypeSize(datatype), comm->nRanks, - rcclParamGinAllReduceForceEnable() == 1, ginAllReduceGinMinBytes(comm)); + const size_t typeSize = ncclTypeSize(datatype); + if (!ginAllReduceSizePolicyEligible(count, typeSize, comm->nRanks, rcclParamGinAllReduceForceEnable() == 1, + ginAllReduceGinMinBytes(comm))) { + return false; + } + const size_t bytes = count * typeSize; + // Above the GinScatter cap, leave the collective to the other selectors. + // The gfx950 two-shot kernel is not the fallback for that range. + if (ginAllReduceIsGfx1250(comm) && bytes > kGinAllReduceGinScatterMaxBytes) return false; + if (!ginAllReduceGinScatterDispatch(ginAllReduceIsGfx1250(comm), bytes)) return true; + if (comm->nRanks <= 0 || count % static_cast(comm->nRanks) != 0) return false; + const int64_t scratch = ncclParamGinAllReduceScratchBytes(); + if (scratch <= 0) return false; + const size_t chunkBytes = (count / static_cast(comm->nRanks)) * typeSize; + return ginAllReduceGinScatterLaunch(bytes, chunkBytes, comm->nRanks, static_cast(scratch)); } bool ncclAllReduceGinSdmaYieldToDda(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count, @@ -390,14 +424,15 @@ ncclResult_t ncclAllReduceGinSdma(const void* sendbuff, void* recvbuff, size_t c ncclResult_t ncclGinAllReduceFinalize(ncclComm* comm) { struct ncclGinAllReduceState* state = &comm->ginAllReduceState; - NCCLCHECK(ginAllReduceFreeScratch(comm)); + ncclResult_t ret = ncclSuccess; + NCCLCHECKIGNORE(ginAllReduceFreeScratch(comm), ret); if (state->intraGpuCtaBar != nullptr) { - NCCLCHECK(ncclCudaFree(state->intraGpuCtaBar, comm->memManager)); + NCCLCHECKIGNORE(ncclCudaFree(state->intraGpuCtaBar, comm->memManager), ret); state->intraGpuCtaBar = nullptr; } if (state->initialized) { - NCCLCHECK(ncclDevCommDestroy(comm, &state->devComm)); + NCCLCHECKIGNORE(ncclDevCommDestroy(comm, &state->devComm), ret); state->initialized = false; } - return ncclSuccess; + return ret; } diff --git a/projects/rccl/src/algorithms/gin/sdma/gin_scatter_reduce.h b/projects/rccl/src/algorithms/gin/sdma/gin_scatter_reduce.h new file mode 100644 index 000000000000..2ceb9f958c61 --- /dev/null +++ b/projects/rccl/src/algorithms/gin/sdma/gin_scatter_reduce.h @@ -0,0 +1,112 @@ +/************************************************************************* + * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. + * + * CE-style local reduce for GinScatter. Self is read from the send slice; + * every other rank is read from that rank's incoming scratch column. + * See LICENSE.txt for license information. + ******************************************************************************/ + +#pragma once + +#include "algorithms/dda/device/CollCommon.h" + +namespace gin::sdma { + +// 16B window accesses through the global aperture. +// +// ncclGetLsaPointer hands back a generic (flat) pointer built by bit-twiddling a base loaded from +// the window struct, so nothing downstream can prove it is device memory and the vector accesses +// compile to flat_load_dwordx4 / flat_store_dwordx4. Peer windows live in the global aperture, so +// casting to an address_space(1) pointer leaves the semantics untouched (plain, non-atomic, ordered +// by the surrounding LSA barriers) while emitting global_load_dwordx4 / global_store_dwordx4 — +// exactly what rccl_ptr.h prescribes for hot paths. If an ISA dump of this kernel mentions "flat", +// one of these casts was lost. +__device__ __forceinline__ uint4 lsaLoadVec(const char* p) { + union { + v4u vec; + uint4 val; + } u; + u.vec = *(v4u_gptr)(const_cast(p)); + return u.val; +} + +__device__ __forceinline__ void lsaStoreVec(char* p, uint4 val) { + union { + v4u vec; + uint4 val; + } u; + u.val = val; + *(v4u_gptr)(p) = u.vec; +} + +// CE ncclCeLocalReduceKernelVec: 16B vectors, GpuUnroll=4, rank loop 0..nRanks-1 +// into recvbuff + rank*shard. Self is not written to scratch (GIN-put skips it), +// so rank == myRank loads from send instead of the incoming slot. +template +__device__ __forceinline__ const char* ginScatterReduceSrc(int peer, int rank, const char* selfSend, + const char* incomingBase, size_t incomingStride, size_t elem, + size_t elemStart, bool staged) { + if (peer == rank) { + return selfSend + elem * sizeof(T); + } + const size_t srcOff = staged ? (elem - elemStart) * sizeof(T) : elem * sizeof(T); + return incomingBase + static_cast(peer) * incomingStride + srcOff; +} + +template +__device__ __forceinline__ void ginScatterReduceChunk(const char* selfSend, const char* incomingBase, char* reducedOut, + int rank, int nRanks, size_t incomingStride, size_t elemStart, + size_t elemEnd, bool staged) { + constexpr int W = static_cast(sizeof(uint4) / sizeof(T)); + constexpr int U = 4; + const size_t nVec = (elemEnd - elemStart) / static_cast(W); + const size_t tid = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + const size_t stride = static_cast(blockDim.x) * gridDim.x; + // Complete U-wide spans only. The tail used to restart at the last partial + // span, so threads whose unrolled iteration already covered that span reduced + // those vectors twice. + const size_t span = stride * static_cast(U); + const size_t nComplete = span == 0 ? 0 : (nVec / span) * span; + + for (size_t vi = tid; vi + static_cast(U - 1) * stride < nComplete; vi += span) { + uint4 acc[U]; + size_t elem[U]; +#pragma unroll + for (int i = 0; i < U; ++i) { + elem[i] = elemStart + (vi + static_cast(i) * stride) * static_cast(W); + acc[i] = lsaLoadVec(ginScatterReduceSrc(0, rank, selfSend, incomingBase, incomingStride, elem[i], elemStart, + staged)); + } +#pragma clang loop unroll(disable) vectorize(disable) + for (int r = 1; r < nRanks; ++r) { + uint4 tmp[U]; +#pragma unroll + for (int i = 0; i < U; ++i) { + tmp[i] = lsaLoadVec(ginScatterReduceSrc(r, rank, selfSend, incomingBase, incomingStride, elem[i], elemStart, + staged)); + } +#pragma unroll + for (int i = 0; i < U; ++i) { + acc[i] = dda::common::vecElementAdd(acc[i], tmp[i]); + } + } +#pragma unroll + for (int i = 0; i < U; ++i) { + lsaStoreVec(reducedOut + elem[i] * sizeof(T), acc[i]); + } + } + + for (size_t vIdx = nComplete + tid; vIdx < nVec; vIdx += stride) { + const size_t e = elemStart + vIdx * static_cast(W); + uint4 acc = + lsaLoadVec(ginScatterReduceSrc(0, rank, selfSend, incomingBase, incomingStride, e, elemStart, staged)); +#pragma clang loop unroll(disable) vectorize(disable) + for (int r = 1; r < nRanks; ++r) { + acc = dda::common::vecElementAdd( + acc, lsaLoadVec(ginScatterReduceSrc(r, rank, selfSend, incomingBase, incomingStride, e, elemStart, staged))); + } + lsaStoreVec(reducedOut + e * sizeof(T), acc); + } +} + +} // namespace gin::sdma diff --git a/projects/rccl/src/gin/README.md b/projects/rccl/src/gin/README.md index 1300a3f89886..69046b6bdd76 100644 --- a/projects/rccl/src/gin/README.md +++ b/projects/rccl/src/gin/README.md @@ -53,7 +53,9 @@ for the alltoall\_wg offload path. Does not enable GIN plugins. | `NCCL_GIN_ANVIL_SDMA_THRESHOLD` | `128` | Minimum message size (bytes) to use SDMA; smaller messages fall back to the IB proxy | | `NCCL_GIN_ANVIL_SDMA_FUSED_SIGNAL` | `0` | Enable fused signal mode for SDMA (experimental) | | `NCCL_GIN_ANVIL_SDMA_CONN_CHECK` | `1` | Validate LSA signal connectivity on first bind (LSA team only; skipped under `NCCL_GIN_CONNECTION_RAIL`); set exactly `0` to bypass | -| `RCCL_GIN_ALLREDUCE_FORCE_ENABLE` | `0` | GIN AllReduce is used only for messages >= 256 MiB by default (smaller sizes use DDA). Set to `1` to also use GIN AllReduce for smaller messages (LSA one-shot / LSA two-shot). | +| `RCCL_GIN_ALLREDUCE_FORCE_ENABLE` | `0` | GIN AllReduce is used only for messages >= 256 MiB by default on gfx950 (smaller sizes use DDA). Set to `1` to also use GIN AllReduce for smaller messages (LSA one-shot / LSA two-shot). | +| `NCCL_GIN_ALLREDUCE_GFX1250_ENABLE` | `0` | Opt in to GinScatter AllReduce on gfx1250 for symmetric `ncclSum` fp32/fp16/bf16 messages from 64 MiB through 4 GiB. `0` leaves gfx1250 on the pre-existing selectors. | +| `NCCL_GIN_ALLREDUCE_SCRATCH_BYTES` | `536870912` | GinScatter incoming scratch, in bytes (512 MiB, matching CE). `0` skips the allocation and GinScatter is not selected. | | `NCCL_CUMEM_ENABLE` | `0` | Required: GIN needs `hipMemCreate`-based allocations | | `NCCL_DMABUF_ENABLE` | `0` | Recommended: enables dmabuf-based MR registration | | `NCCL_P2P_DISABLE` | `0` | Set to `1` to force inter-GPU traffic over the network (useful for single-node GIN testing) | diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index 2f918707dfa4..dbbcabea79f1 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -1426,8 +1426,9 @@ ncclResult_t rcclSelectAllReduce(struct ncclComm* comm, const void* sendbuff, vo // registered operands to DDA. // // GIN AllReduce is selected first in this function and requires symmetric - // windows. By default gfx950 claims messages >= 256 MiB and gfx1250 with - // Anvil SDMA claims >= 64 MiB, so DDA must still be allowed for smaller + // windows. By default gfx950 claims messages >= 256 MiB. gfx1250 claims + // [64 MiB, 4 GiB] only when NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1, so DDA must + // still be allowed for smaller // symmetric AllReduces (otherwise they would hit the symmetric kernel // instead of DDA). FORCE_ENABLE=1 keeps the original // !symEligible gate because GIN already returned above for those sizes. diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index fbbd0410b47c..f2b124f4277d 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -653,6 +653,32 @@ if(BUILD_TESTS) add_test(NAME rccl-UnitTestsGinSdmaPolicy COMMAND rccl-UnitTestsGinSdmaPolicy) set_tests_properties(rccl-UnitTestsGinSdmaPolicy PROPERTIES LABELS "unit") + # GinScatter local reduce on one GPU. Runs ginScatterReduceChunk against a fake + # scratch and a host sum, for staged and full-column layouts. No librccl. + # Skips when no GPU is visible. + add_executable(rccl-UnitTestsGinScatterReduceGpu gin/gin_scatter_reduce_gpu_test.cpp) + set_source_files_properties(gin/gin_scatter_reduce_gpu_test.cpp PROPERTIES LANGUAGE CXX) + target_include_directories(rccl-UnitTestsGinScatterReduceGpu PRIVATE + ${PROJECT_SOURCE_DIR}/src/include + ${PROJECT_SOURCE_DIR}/src) + set(_gin_scatter_gpu_offload "") + foreach(_g IN LISTS GPU_TARGETS) + if(_g) + list(APPEND _gin_scatter_gpu_offload "--offload-arch=${_g}") + endif() + endforeach() + if(_gin_scatter_gpu_offload) + target_compile_options(rccl-UnitTestsGinScatterReduceGpu PRIVATE ${_gin_scatter_gpu_offload}) + target_link_options(rccl-UnitTestsGinScatterReduceGpu PRIVATE ${_gin_scatter_gpu_offload}) + endif() + target_link_libraries(rccl-UnitTestsGinScatterReduceGpu PRIVATE + ${GTEST_BOTH_LIBRARIES} hip::host hip::device Threads::Threads) + set_target_properties(rccl-UnitTestsGinScatterReduceGpu PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/test" + CXX_STANDARD 17 CXX_STANDARD_REQUIRED ON) + add_test(NAME rccl-UnitTestsGinScatterReduceGpu COMMAND rccl-UnitTestsGinScatterReduceGpu) + set_tests_properties(rccl-UnitTestsGinScatterReduceGpu PROPERTIES LABELS "unit;gpu") + # GIN Anvil-SDMA hybrid AllGather (-D 3) policy host unit test (no GPU). # Compiles gin_sdma_allgather_policy.h from the sibling rccl-tests project with # GIN_SDMA_HOST_ONLY (no librccl / device codegen needed), validating the exact diff --git a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp index ad63c097da92..7a915c6bc756 100644 --- a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp +++ b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp @@ -33,7 +33,7 @@ TEST(GinAllReducePolicy, ThresholdsMatchDocumentedBands) { EXPECT_EQ(kGinAllReduceMaxRanks, 8); EXPECT_EQ(kGinAllReduceMinPutBytes, 128u); EXPECT_EQ(kGinAllReduceTwoShotScratchBytes, 512ull * kMiB); - EXPECT_EQ(kGinAllReduceGinScatterChunkBytes, 16ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterChunkBytes, 64ull * kMiB); EXPECT_EQ(kGinAllReduceGinScatterNumSlots, 2); } @@ -114,8 +114,8 @@ TEST(GinAllReducePolicy, TwoShotRejectsCountNotDivisibleByRanks) { } TEST(GinAllReducePolicy, Gfx1250ScatterMinBytesFrom64MiB) { - EXPECT_FALSE( - ginAllReduceSizePolicyEligible(countForBytes(64 * kMiB - 16), kFloat, kRanks, false, kGinAllReduceGinScatterMinBytes)); + EXPECT_FALSE(ginAllReduceSizePolicyEligible(countForBytes(64 * kMiB - 16), kFloat, kRanks, false, + kGinAllReduceGinScatterMinBytes)); EXPECT_TRUE( ginAllReduceSizePolicyEligible(countForBytes(64 * kMiB), kFloat, kRanks, false, kGinAllReduceGinScatterMinBytes)); EXPECT_TRUE( @@ -133,6 +133,8 @@ TEST(GinAllReducePolicy, GinScatterDispatchGfx1250From64MiB) { EXPECT_FALSE(ginAllReduceGinScatterDispatch(true, 64ull * kMiB - 16)); EXPECT_TRUE(ginAllReduceGinScatterDispatch(true, 64ull * kMiB)); EXPECT_TRUE(ginAllReduceGinScatterDispatch(true, 256ull * kMiB)); + EXPECT_TRUE(ginAllReduceGinScatterDispatch(true, kGinAllReduceGinScatterMaxBytes)); + EXPECT_FALSE(ginAllReduceGinScatterDispatch(true, kGinAllReduceGinScatterMaxBytes + 16)); EXPECT_FALSE(ginAllReduceGinScatterDispatch(false, 64ull * kMiB)); EXPECT_FALSE(ginAllReduceGinScatterDispatch(false, 256ull * kMiB)); } @@ -148,6 +150,8 @@ TEST(GinAllReducePolicy, GinScatterLaunchFrom64MiB) { EXPECT_TRUE(ginAllReduceGinScatterLaunch(512ull * kMiB, chunk512, kRanks, scratch)); EXPECT_TRUE(ginAllReduceGinScatterLaunch(2ull * 1024ull * kMiB, 256ull * kMiB, kRanks, scratch)); EXPECT_TRUE(ginAllReduceGinScatterLaunch(kGinAllReduceGinScatterMaxBytes, 512ull * kMiB, kRanks, scratch)); + EXPECT_FALSE(ginAllReduceGinScatterLaunch(kGinAllReduceGinScatterMaxBytes + 16, 512ull * kMiB, kRanks, scratch)); + EXPECT_FALSE(ginAllReduceGinScatterLaunch(64ull * kMiB, chunk64, kRanks, 0)); } TEST(GinAllReducePolicy, GinScatterLaunchWhenMessageExceedsScratch) { @@ -173,14 +177,16 @@ TEST(GinAllReducePolicy, GinScatterStagesWhenMessageExceedsScratch) { int nPhases = 0; size_t uniformBytes = 0; size_t lastBytes = 0; - const size_t smallSeq = ginAllReduceGinScatterStagedScratchBytes(kGinAllReduceGinScatterChunkBytes, kRanks, 1); - ginAllReduceGinScatterChunkPlan(32ull * kMiB, kRanks, smallSeq, &nChunks, &uniformBytes, &lastBytes, &nPhases); + const int seqRanks = 4; + const size_t smallSeq = ginAllReduceGinScatterStagedScratchBytes(kGinAllReduceGinScatterChunkBytes, seqRanks, 1); + ginAllReduceGinScatterChunkPlan(96ull * kMiB, seqRanks, smallSeq, &nChunks, &uniformBytes, &lastBytes, &nPhases); EXPECT_EQ(nChunks, 2); EXPECT_EQ(nPhases, 1); EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); + EXPECT_EQ(lastBytes, 32ull * kMiB); const size_t smallPipe = ginAllReduceGinScatterStagedScratchBytes(kGinAllReduceGinScatterChunkBytes, kRanks, 2); ginAllReduceGinScatterChunkPlan(512ull * kMiB, kRanks, smallPipe, &nChunks, &uniformBytes, &lastBytes, &nPhases); - EXPECT_EQ(nChunks, 32); + EXPECT_EQ(nChunks, 8); EXPECT_EQ(nPhases, 2); EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); } @@ -205,14 +211,14 @@ TEST(GinAllReducePolicy, GinScatterNoPipelineBelow512MiB) { EXPECT_EQ(nPhases, 1); } -TEST(GinAllReducePolicy, GinScatterChunksOverlapAt512MiB) { +TEST(GinAllReducePolicy, GinScatterOneChunkAt512MiB) { int nChunks = 0; int nPhases = 0; size_t uniformBytes = 0; size_t lastBytes = 0; ginAllReduceGinScatterChunkPlan(64ull * kMiB, kRanks, kGinAllReduceTwoShotScratchBytes, &nChunks, &uniformBytes, &lastBytes, &nPhases); - EXPECT_EQ(nChunks, 4); + EXPECT_EQ(nChunks, 1); EXPECT_EQ(nPhases, 1); EXPECT_EQ(uniformBytes, kGinAllReduceGinScatterChunkBytes); EXPECT_EQ(lastBytes, kGinAllReduceGinScatterChunkBytes); @@ -268,10 +274,10 @@ TEST(GinAllReducePolicy, GinScatterChunksTilePerRankBytes) { } TEST(GinAllReducePolicy, GinScatterScratchIsIncomingSlots) { - EXPECT_EQ(ginAllReduceGinScatterScratchBytes(32ull * kMiB, kRanks), 8ull * 32ull * kMiB); + EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(32ull * kMiB, kRanks), 8ull * 32ull * kMiB); EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(8ull * kMiB, kRanks, 2), 16ull * 8ull * kMiB); - EXPECT_EQ(ginAllReduceGinScatterScratchBytes(16, 1), 16u); - EXPECT_EQ(ginAllReduceGinScatterScratchBytes(16, 0), 0u); + EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(16, 1), 16u); + EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(16, 0), 0u); EXPECT_EQ(ginAllReduceGinScatterFitSlot(kGinAllReduceTwoShotScratchBytes, kRanks, 1), 64ull * kMiB); EXPECT_EQ(ginAllReduceGinScatterFitSlot(kGinAllReduceTwoShotScratchBytes, kRanks, 2), 32ull * kMiB); } diff --git a/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp new file mode 100644 index 000000000000..d7749a6bd5b5 --- /dev/null +++ b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp @@ -0,0 +1,128 @@ +/************************************************************************* + * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Single-GPU check of ginScatterReduceChunk. Fills a fake incoming scratch +// with known per-peer patterns and compares the reduced slice to a host sum. +// Covers staged and full-column layouts, including a chunk large enough that +// the 56x512 unroll used to reduce the tail twice. No librccl and no GIN +// communicator. Skips when no GPU is visible. + +#include + +#include +#include +#include + +#include + +#include "algorithms/gin/gin_all_reduce_policy.h" +#include "algorithms/gin/sdma/gin_scatter_reduce.h" + +namespace { + +#define HIP_OK(cmd) \ + do { \ + hipError_t _e = (cmd); \ + ASSERT_EQ(_e, hipSuccess) << "HIP error " << (int)_e << " (" << hipGetErrorString(_e) \ + << ") at " << __FILE__ << ":" << __LINE__; \ + } while (0) + +bool gpuAvailable() { + int n = 0; + hipError_t e = hipGetDeviceCount(&n); + return e == hipSuccess && n > 0; +} + +__global__ void ginScatterReduceChunkKernel(const float* selfSend, const float* incoming, float* reducedOut, int rank, + int nRanks, size_t incomingStride, size_t elemStart, size_t elemEnd, + bool staged) { + gin::sdma::ginScatterReduceChunk(reinterpret_cast(selfSend), + reinterpret_cast(incoming), + reinterpret_cast(reducedOut), rank, nRanks, incomingStride, elemStart, + elemEnd, staged); +} + +float hostSum(const std::vector& selfSend, const std::vector& incoming, int rank, int nRanks, + size_t strideElems, size_t elem, size_t elemStart, bool staged) { + float sum = 0.f; + for (int peer = 0; peer < nRanks; ++peer) { + if (peer == rank) { + sum += selfSend[elem]; + } else if (staged) { + sum += incoming[static_cast(peer) * strideElems + (elem - elemStart)]; + } else { + sum += incoming[static_cast(peer) * strideElems + elem]; + } + } + return sum; +} + +void runReduceCase(bool staged, size_t elemStart, size_t nVec, int rank, int nRanks) { + constexpr int W = static_cast(sizeof(uint4) / sizeof(float)); + const size_t chunkElems = nVec * static_cast(W); + const size_t elemEnd = elemStart + chunkElems; + const size_t strideElems = staged ? chunkElems : elemEnd; + const size_t incomingStride = strideElems * sizeof(float); + + std::vector selfSend(elemEnd, 0.f); + std::vector incoming(static_cast(nRanks) * strideElems, -1.f); + for (size_t elem = elemStart; elem < elemEnd; ++elem) { + selfSend[elem] = static_cast(rank * 1000 + static_cast(elem % 64)); + for (int peer = 0; peer < nRanks; ++peer) { + if (peer == rank) continue; + const size_t src = staged ? static_cast(peer) * strideElems + (elem - elemStart) + : static_cast(peer) * strideElems + elem; + incoming[src] = static_cast(peer * 1000 + static_cast(elem % 64)); + } + } + + float *dSelf = nullptr, *dIncoming = nullptr, *dOut = nullptr; + HIP_OK(hipMalloc(&dSelf, selfSend.size() * sizeof(float))); + HIP_OK(hipMalloc(&dIncoming, incoming.size() * sizeof(float))); + HIP_OK(hipMalloc(&dOut, elemEnd * sizeof(float))); + HIP_OK(hipMemcpy(dSelf, selfSend.data(), selfSend.size() * sizeof(float), hipMemcpyHostToDevice)); + HIP_OK(hipMemcpy(dIncoming, incoming.data(), incoming.size() * sizeof(float), hipMemcpyHostToDevice)); + HIP_OK(hipMemset(dOut, 0, elemEnd * sizeof(float))); + + ginScatterReduceChunkKernel<<>>( + dSelf, dIncoming, dOut, rank, nRanks, incomingStride, elemStart, elemEnd, staged); + HIP_OK(hipGetLastError()); + HIP_OK(hipDeviceSynchronize()); + + std::vector got(elemEnd, 1.f); + HIP_OK(hipMemcpy(got.data(), dOut, elemEnd * sizeof(float), hipMemcpyDeviceToHost)); + + size_t wrong = 0; + for (size_t elem = 0; elem < elemStart; ++elem) { + if (got[elem] != 0.f) ++wrong; + } + for (size_t elem = elemStart; elem < elemEnd; ++elem) { + const float exp = hostSum(selfSend, incoming, rank, nRanks, strideElems, elem, elemStart, staged); + if (got[elem] != exp) ++wrong; + } + EXPECT_EQ(wrong, 0u) << "staged=" << staged << " elemStart=" << elemStart << " nVec=" << nVec << " rank=" << rank + << " nRanks=" << nRanks; + + HIP_OK(hipFree(dSelf)); + HIP_OK(hipFree(dIncoming)); + HIP_OK(hipFree(dOut)); +} + +} // namespace + +TEST(GinScatterReduceGpu, ChunkMatchesHostSum) { + if (!gpuAvailable()) GTEST_SKIP() << "no visible GPU"; + + const size_t stride = static_cast(kGinAllReduceLsaCtas) * static_cast(kGinAllReduceLsaThreadsPerCta); + const size_t span = stride * 4u; + // One full unroll span plus a tail. The old tail restarted inside that span. + const size_t overlapping = span + 1024u; + + runReduceCase(/*staged=*/false, /*elemStart=*/0, overlapping, /*rank=*/1, /*nRanks=*/4); + runReduceCase(/*staged=*/true, /*elemStart=*/256, overlapping, /*rank=*/0, /*nRanks=*/4); + runReduceCase(/*staged=*/false, /*elemStart=*/0, /*nVec=*/100, /*rank=*/0, /*nRanks=*/2); + runReduceCase(/*staged=*/true, /*elemStart=*/64, /*nVec=*/100, /*rank=*/1, /*nRanks=*/3); +} From 126f5e01ec29a576ced72359ae9506fffecb95dd Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Thu, 22 Oct 2026 09:50:57 -0500 Subject: [PATCH 3/7] address reviews --- projects/rccl/cmake/DeviceLinker.cmake | 112 +++++++++--------- .../algorithms/gin/gin_all_reduce_policy.h | 19 +++ .../algorithms/gin/sdma/all_reduce_gin_sdma.h | 35 ++---- .../test/gin/gin_all_reduce_policy_test.cpp | 38 ++++++ .../test/gin/gin_scatter_reduce_gpu_test.cpp | 10 +- 5 files changed, 134 insertions(+), 80 deletions(-) diff --git a/projects/rccl/cmake/DeviceLinker.cmake b/projects/rccl/cmake/DeviceLinker.cmake index 0d46146e9aaf..5e5f12e79020 100644 --- a/projects/rccl/cmake/DeviceLinker.cmake +++ b/projects/rccl/cmake/DeviceLinker.cmake @@ -915,56 +915,6 @@ if(ENABLE_ROCSHMEM_GIN) ) endif() -# =========================================================================== -# rocshmem QP device bitcode for GIN GDA puts (QueuePair::put_nbi, etc.). -# Used by GIN symmetric kernels. The installed per-arch .bc files are -# arch-optimized and cannot be used with -mlink-builtin-bitcode across archs, -# so this is unoptimized IR + constmem stubs. -# =========================================================================== -set(GIN_ROCSHMEM_QP_BC_FLAG "") -set(GIN_ROCSHMEM_QP_DEPS "") -if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR AND DL_GPU_TARGETS AND GENERATE_SYM_KERNELS) - list(GET DL_GPU_TARGETS 0 _bc_arch) - set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc") - find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED) - find_program(_llvm_dis llvm-dis HINTS ${ROCM_PATH}/llvm/bin REQUIRED) - find_program(_llvm_as llvm-as HINTS ${ROCM_PATH}/llvm/bin REQUIRED) - - set(_cm_src "${CMAKE_SOURCE_DIR}/src/gin/gin_rocshmem_constmem.hip") - set(_cm_bc "${DEVICE_BUILD_DIR}/gin_rocshmem_constmem.bc") - set(_qp_raw "${DEVICE_BUILD_DIR}/rocshmem_qp_raw.bc") - - add_custom_command( - OUTPUT ${_cm_bc} - COMMAND ${DL_CLANG} - -x hip --cuda-device-only --offload-arch=${_bc_arch} - -emit-llvm -Xclang -disable-llvm-passes - -std=c++17 -fPIC - -I${ROCSHMEM_SOURCE_DIR}/src - -I${ROCSHMEM_SOURCE_DIR}/include - -c -o ${_cm_bc} ${_cm_src} - DEPENDS ${_cm_src} - COMMENT "DL: compiling rocshmem constmem stubs to device bitcode" - VERBATIM) - - add_custom_command( - OUTPUT ${_qp_bc} - COMMAND ${_llvm_link} ${_cm_bc} -o ${_qp_raw} - COMMAND ${_llvm_dis} -o ${_qp_raw}.ll ${_qp_raw} - COMMAND grep -v -E "@llvm[.]compiler[.]used|@__hip_cuid_" - ${_qp_raw}.ll > ${_qp_raw}.clean.ll - COMMAND ${_llvm_as} ${_qp_raw}.clean.ll -o ${_qp_bc} - DEPENDS ${_cm_bc} - COMMENT "DL: linking rocshmem QP device bitcode (with constmem, stripped)" - VERBATIM) - add_custom_target(dl_rocshmem_qp_bc DEPENDS ${_qp_bc}) - if(TARGET rocshmem_static) - add_dependencies(dl_rocshmem_qp_bc rocshmem_static) - endif() - set(GIN_ROCSHMEM_QP_BC_FLAG -Xclang -mlink-builtin-bitcode -Xclang ${_qp_bc}) - set(GIN_ROCSHMEM_QP_DEPS dl_rocshmem_qp_bc) -endif() - # =========================================================================== # gin_all_reduce_sdma.cu.cpp: GIN-SDMA allreduce kernel. # @@ -1351,9 +1301,65 @@ endforeach() # =========================================================================== set(SYM_FAT_OBJS "") if(GENERATE_SYM_KERNELS) - # GIN symmetric kernels reuse the QP bitcode generated above. - set(_sym_rocshmem_bc_flag "${GIN_ROCSHMEM_QP_BC_FLAG}") - set(_sym_rocshmem_deps "${GIN_ROCSHMEM_QP_DEPS}") + # When ENABLE_ROCSHMEM_GIN is set, GIN device templates reference rocshmem + # device symbols (QueuePair::put_nbi, atomic_add, etc.). The installed per-arch + # .bc files are arch-optimized (opt -mcpu=) and can't be used with + # -mlink-builtin-bitcode across archs. Instead, llvm-link the pre-opt + # individual source .bc files (arch-agnostic) into a minimal QP-only bitcode. + # GinScatter does not use this bitcode, so it stays inside the symmetric-kernel build. + set(_sym_rocshmem_bc_flag "") + set(_sym_rocshmem_deps "") + if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR) + # Pick the first arch's pre-opt bitcode dir (all archs produce identical + # unoptimized IR since -Xclang -disable-llvm-passes is used). + list(GET DL_GPU_TARGETS 0 _bc_arch) + set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc") + find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED) + find_program(_llvm_dis llvm-dis HINTS ${ROCM_PATH}/llvm/bin REQUIRED) + find_program(_llvm_as llvm-as HINTS ${ROCM_PATH}/llvm/bin REQUIRED) + + # Pipeline: + # 1. Compile gin_rocshmem_constmem.hip → device-only .bc (provides + # rocshmem::constmem and rocshmem::logd_constants definitions that + # queue_pair.bc references as external) + # 2. llvm-link QP .bc files + constmem .bc into one module + # 3. Strip @llvm.compiler.used and @__hip_cuid_ via text round-trip + # (these AMDGCN addrspace(1) appending globals clash with the + # host-side addrspace(0) equivalents in fat-object compilation) + set(_cm_src "${CMAKE_SOURCE_DIR}/src/gin/gin_rocshmem_constmem.hip") + set(_cm_bc "${DEVICE_BUILD_DIR}/gin_rocshmem_constmem.bc") + set(_qp_raw "${DEVICE_BUILD_DIR}/rocshmem_qp_raw.bc") + + add_custom_command( + OUTPUT ${_cm_bc} + COMMAND ${DL_CLANG} + -x hip --cuda-device-only --offload-arch=${_bc_arch} + -emit-llvm -Xclang -disable-llvm-passes + -std=c++17 -fPIC + -I${ROCSHMEM_SOURCE_DIR}/src + -I${ROCSHMEM_SOURCE_DIR}/include + -c -o ${_cm_bc} ${_cm_src} + DEPENDS ${_cm_src} + COMMENT "DL: compiling rocshmem constmem stubs to device bitcode" + VERBATIM) + + add_custom_command( + OUTPUT ${_qp_bc} + COMMAND ${_llvm_link} ${_cm_bc} -o ${_qp_raw} + COMMAND ${_llvm_dis} -o ${_qp_raw}.ll ${_qp_raw} + COMMAND grep -v -E "@llvm[.]compiler[.]used|@__hip_cuid_" + ${_qp_raw}.ll > ${_qp_raw}.clean.ll + COMMAND ${_llvm_as} ${_qp_raw}.clean.ll -o ${_qp_bc} + DEPENDS ${_cm_bc} + COMMENT "DL: linking rocshmem QP device bitcode (with constmem, stripped)" + VERBATIM) + add_custom_target(dl_rocshmem_qp_bc DEPENDS ${_qp_bc}) + if(TARGET rocshmem_static) + add_dependencies(dl_rocshmem_qp_bc rocshmem_static) + endif() + set(_sym_rocshmem_bc_flag -Xclang -mlink-builtin-bitcode -Xclang ${_qp_bc}) + set(_sym_rocshmem_deps dl_rocshmem_qp_bc) + endif() file(GLOB _sym_srcs CONFIGURE_DEPENDS "${HIPIFY_DIR}/gensrc/symmetric/*.cpp") foreach(_sym_src IN LISTS _sym_srcs) get_filename_component(_sym_name "${_sym_src}" NAME_WE) diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h index cd0f084f8027..1af758554665 100644 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h @@ -151,6 +151,25 @@ GIN_AR_HD size_t ginAllReduceGinScatterChunkSize(int chunk, int nChunks, size_t return (chunk == nChunks - 1) ? lastBytes : uniformBytes; } +// Staged scratch is [slot][rank] columns of slotBytes. Unstaged is one column +// per rank of perRankBytes, plus the chunk offset within that column. +GIN_AR_HD size_t ginScatterPutDestOff(size_t scratchOff, int rank, int nRanks, size_t perRankBytes, size_t slotBytes, + int slot, size_t chunkOff, bool staged) { + if (!staged) { + return scratchOff + static_cast(rank) * perRankBytes + chunkOff; + } + return scratchOff + static_cast(slot) * static_cast(nRanks) * slotBytes + + static_cast(rank) * slotBytes; +} + +GIN_AR_HD const char* ginScatterIncomingBase(const char* scratchBase, int nRanks, size_t slotBytes, int slot, + bool staged) { + if (!staged) { + return scratchBase; + } + return scratchBase + static_cast(slot) * static_cast(nRanks) * slotBytes; +} + #undef GIN_AR_HD // Two-shot kernels require a whole number of elements per rank and a 16-byte diff --git a/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h b/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h index 69bec2a3b2de..08dad765a83e 100755 --- a/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h +++ b/projects/rccl/src/algorithms/gin/sdma/all_reduce_gin_sdma.h @@ -338,23 +338,6 @@ __device__ __forceinline__ void ginScatterPutRemoteChunk(ncclGin& gin, ncclTeam } } -__device__ __forceinline__ size_t ginScatterPutDestOff(size_t scratchOff, int rank, int nRanks, size_t perRankBytes, - size_t slotBytes, int slot, size_t chunkOff, bool staged) { - if (!staged) { - return scratchOff + static_cast(rank) * perRankBytes + chunkOff; - } - return scratchOff + static_cast(slot) * static_cast(nRanks) * slotBytes + - static_cast(rank) * slotBytes; -} - -__device__ __forceinline__ const char* ginScatterIncomingBase(const char* scratchBase, int nRanks, size_t slotBytes, - int slot, bool staged) { - if (!staged) { - return scratchBase; - } - return scratchBase + static_cast(slot) * static_cast(nRanks) * slotBytes; -} - __device__ __forceinline__ void ginScatterIssuePut(ncclGin& gin, ncclTeam world, int rank, int nRanks, ncclWindow_t scratchWin, size_t scratchOff, ncclWindow_t sendWin, size_t sendOff, size_t perRankBytes, size_t slotBytes, @@ -378,9 +361,11 @@ __device__ __forceinline__ void ginScatterWaitThread0(ncclGin& gin, unsigned sig } // Thread 0 waits until every other rank's own signal has advanced by `add`. -// `baseline` is this thread's snapshot and is only valid on thread 0. -__device__ __forceinline__ void ginScatterWaitPeersThread0(ncclGin& gin, unsigned base, const uint64_t* baseline, - int rank, int nRanks, uint64_t add) { +// `baseline` is the CTA-shared snapshot. A reference keeps HIP's shared address +// space; a plain uint64_t* cannot bind a __shared__ array. +template +__device__ __forceinline__ void ginScatterWaitPeersThread0(ncclGin& gin, unsigned base, Baseline& baseline, int rank, + int nRanks, uint64_t add) { if (threadIdx.x != 0) return; for (int peer = 0; peer < nRanks; ++peer) { if (peer == rank) continue; @@ -427,6 +412,12 @@ __launch_bounds__(512) size_t scratchBytes, uint32_t* intraGpuCtaBar) { constexpr int ginContext = 0; + // Thread 0 of CTA 0 snapshots these. One shared copy per CTA, not a + // private array on every thread. + __shared__ uint64_t scatterBase[kGinAllReduceMaxRanks]; + __shared__ uint64_t doneBase[kGinAllReduceMaxRanks]; + __shared__ uint64_t agValue; + // gin/ginBar must outlive the first put; only CTA 0 snapshots signals and world-syncs. // Pipeline steps use per-peer waitSignal / WeakSignalInc, not a shared counter. ncclGin gin{devComm, ginContext}; @@ -451,10 +442,6 @@ __launch_bounds__(512) const size_t incomingStride = staged ? slotBytes : perRankBytes; const bool overlap = (nChunks > 1) && !(staged && nPhases <= 1); - uint64_t scatterBase[kGinAllReduceMaxRanks] = {}; - uint64_t doneBase[kGinAllReduceMaxRanks] = {}; - uint64_t agValue = 0; - // One world acquire so every rank snapshots per-peer baselines before any put. // Later pipeline steps are per-peer waitSignal doorbells, not ginBar.sync. if (blockIdx.x == 0) { diff --git a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp index 7a915c6bc756..d757fbf0cf39 100644 --- a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp +++ b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp @@ -11,6 +11,7 @@ #include #include +#include #include "algorithms/gin/gin_all_reduce_policy.h" @@ -273,6 +274,43 @@ TEST(GinAllReducePolicy, GinScatterChunksTilePerRankBytes) { EXPECT_EQ(ginAllReduceGinScatterChunkSize(nChunks - 1, nChunks, uniformBytes, lastBytes), lastBytes); } +TEST(GinAllReducePolicy, GinScatterChunksTilePerRankBytesUnstaged) { + const int nRanks = 4; + const size_t perRank = 128ull * kMiB; + const size_t scratch = 512ull * kMiB; + EXPECT_FALSE(ginAllReduceGinScatterStaged(perRank, nRanks, scratch)); + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + ginAllReduceGinScatterChunkPlan(perRank, nRanks, scratch, &nChunks, &uniformBytes, &lastBytes, &nPhases); + EXPECT_EQ(nChunks, 2); + EXPECT_EQ(nPhases, 1); + EXPECT_EQ(uniformBytes, 64ull * kMiB); + EXPECT_EQ(lastBytes, 64ull * kMiB); + const bool overlap = (nChunks > 1) && !(ginAllReduceGinScatterStaged(perRank, nRanks, scratch) && nPhases <= 1); + EXPECT_TRUE(overlap); +} + +TEST(GinAllReducePolicy, GinScatterScratchLayout) { + constexpr int nRanks = 4; + constexpr int rank = 2; + constexpr int slot = 1; + constexpr size_t slotBytes = 32; + constexpr size_t scratchOff = 64; + constexpr size_t perRank = 1000; + constexpr size_t chunkOff = 16; + EXPECT_EQ(ginScatterPutDestOff(scratchOff, rank, nRanks, perRank, slotBytes, slot, chunkOff, false), + scratchOff + static_cast(rank) * perRank + chunkOff); + EXPECT_EQ(ginScatterPutDestOff(scratchOff, rank, nRanks, perRank, slotBytes, slot, chunkOff, true), + scratchOff + static_cast(slot) * nRanks * slotBytes + static_cast(rank) * slotBytes); + const auto base = reinterpret_cast(static_cast(0x10000)); + EXPECT_EQ(ginScatterIncomingBase(base, nRanks, slotBytes, slot, false), base); + const char* staged = ginScatterIncomingBase(base, nRanks, slotBytes, slot, true); + EXPECT_EQ(static_cast(staged), + static_cast(base) + static_cast(slot) * nRanks * slotBytes); +} + TEST(GinAllReducePolicy, GinScatterScratchIsIncomingSlots) { EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(32ull * kMiB, kRanks), 8ull * 32ull * kMiB); EXPECT_EQ(ginAllReduceGinScatterStagedScratchBytes(8ull * kMiB, kRanks, 2), 16ull * 8ull * kMiB); diff --git a/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp index d7749a6bd5b5..37ccf09d8d71 100644 --- a/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp +++ b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp @@ -60,11 +60,12 @@ float hostSum(const std::vector& selfSend, const std::vector& inco return sum; } -void runReduceCase(bool staged, size_t elemStart, size_t nVec, int rank, int nRanks) { +void runReduceCase(bool staged, size_t elemStart, size_t nVec, int rank, int nRanks, size_t extraStrideElems = 0) { constexpr int W = static_cast(sizeof(uint4) / sizeof(float)); const size_t chunkElems = nVec * static_cast(W); const size_t elemEnd = elemStart + chunkElems; - const size_t strideElems = staged ? chunkElems : elemEnd; + // Production passes slotBytes or perRankBytes, which exceed the last chunk. + const size_t strideElems = (staged ? chunkElems : elemEnd) + extraStrideElems; const size_t incomingStride = strideElems * sizeof(float); std::vector selfSend(elemEnd, 0.f); @@ -104,7 +105,7 @@ void runReduceCase(bool staged, size_t elemStart, size_t nVec, int rank, int nRa if (got[elem] != exp) ++wrong; } EXPECT_EQ(wrong, 0u) << "staged=" << staged << " elemStart=" << elemStart << " nVec=" << nVec << " rank=" << rank - << " nRanks=" << nRanks; + << " nRanks=" << nRanks << " strideElems=" << strideElems; HIP_OK(hipFree(dSelf)); HIP_OK(hipFree(dIncoming)); @@ -125,4 +126,7 @@ TEST(GinScatterReduceGpu, ChunkMatchesHostSum) { runReduceCase(/*staged=*/true, /*elemStart=*/256, overlapping, /*rank=*/0, /*nRanks=*/4); runReduceCase(/*staged=*/false, /*elemStart=*/0, /*nVec=*/100, /*rank=*/0, /*nRanks=*/2); runReduceCase(/*staged=*/true, /*elemStart=*/64, /*nVec=*/100, /*rank=*/1, /*nRanks=*/3); + // Stride strictly larger than the chunk, matching a last-chunk slot or column. + runReduceCase(/*staged=*/true, /*elemStart=*/256, overlapping, /*rank=*/0, /*nRanks=*/4, /*extraStrideElems=*/64); + runReduceCase(/*staged=*/false, /*elemStart=*/0, /*nVec=*/100, /*rank=*/1, /*nRanks=*/4, /*extraStrideElems=*/32); } From f1e757ab8fb11293389621d5842a782342680770 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Thu, 22 Oct 2026 16:54:28 -0500 Subject: [PATCH 4/7] fix gin allreduce test --- projects/rccl/test/gin/gin_all_reduce_policy_test.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp index d757fbf0cf39..5db7e27ff307 100644 --- a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp +++ b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp @@ -307,8 +307,8 @@ TEST(GinAllReducePolicy, GinScatterScratchLayout) { const auto base = reinterpret_cast(static_cast(0x10000)); EXPECT_EQ(ginScatterIncomingBase(base, nRanks, slotBytes, slot, false), base); const char* staged = ginScatterIncomingBase(base, nRanks, slotBytes, slot, true); - EXPECT_EQ(static_cast(staged), - static_cast(base) + static_cast(slot) * nRanks * slotBytes); + EXPECT_EQ(reinterpret_cast(staged), + reinterpret_cast(base) + static_cast(slot) * nRanks * slotBytes); } TEST(GinAllReducePolicy, GinScatterScratchIsIncomingSlots) { From dd8bbd47b096dd91d170cb19e35a0cea71d068f1 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Mon, 26 Oct 2026 18:49:16 -0500 Subject: [PATCH 5/7] address reviews --- .../rccl/docs/api-reference/env-variables.rst | 2 + .../rccl/src/algorithms/gin/gin_all_reduce.h | 9 +- .../algorithms/gin/gin_all_reduce_policy.h | 3 + .../gin/sdma/gin_all_reduce_sdma.cu | 57 +- projects/rccl/src/rccl_wrap.cc | 9 +- projects/rccl/test/CMakeLists.txt | 5 +- .../test/gin/gin_all_reduce_policy_test.cpp | 5 + .../test/gin/gin_scatter_reduce_gpu_test.cpp | 149 +++- .../configs/mi300_mellanox_ib_func.json | 21 + .../configs/mi300x_mellanox_ib.json | 21 + .../configs/mi355x_thor2_roce.json | 21 + .../configs/mi355x_thor2_roce_func.json | 21 + .../test_runner/configs/mi455_ainic_roce.json | 21 + .../configs/mi455_ainic_roce_func.json | 21 + runtimes/ddi/rocddi/src/topology/snapshot.rs | 682 ------------------ 15 files changed, 332 insertions(+), 715 deletions(-) delete mode 100644 runtimes/ddi/rocddi/src/topology/snapshot.rs diff --git a/projects/rccl/docs/api-reference/env-variables.rst b/projects/rccl/docs/api-reference/env-variables.rst index 492917b79c0f..b6b6166cfdce 100644 --- a/projects/rccl/docs/api-reference/env-variables.rst +++ b/projects/rccl/docs/api-reference/env-variables.rst @@ -695,6 +695,8 @@ in the following table. - | ``0``: gfx1250 keeps the pre-existing AllReduce selectors (default) | ``1``: symmetric ``ncclSum`` fp32/fp16/bf16 messages from 64 MiB through 4 GiB use GinScatter when scratch can hold them + | Eligible sums below 64 MiB yield to DDA. Sums below 8 MiB already + used DDA. From 8 MiB to 64 MiB, DDA LL/LL128 can take those sums * - | ``NCCL_GIN_ALLREDUCE_SCRATCH_BYTES`` | GinScatter incoming scratch size, in bytes. diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce.h b/projects/rccl/src/algorithms/gin/gin_all_reduce.h index 28373f4a54d2..5adeefaf1b5a 100755 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce.h @@ -4,7 +4,8 @@ * Host entry points for the GIN-SDMA AllReduce path selected by rcclSelectAllReduce * when symmetric windows are used. By default gfx950 takes GIN two-shot for * messages >= 256 MiB. gfx1250 GinScatter (64 MiB to 4 GiB) is opt-in via - * NCCL_GIN_ALLREDUCE_GFX1250_ENABLE. Smaller messages use DDA AllReduce. + * NCCL_GIN_ALLREDUCE_GFX1250_ENABLE. Smaller messages yield to DDA. On gfx1250 + * that includes sums from 8 MiB up to 64 MiB, which DDA LL/LL128 can claim. * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables LSA one-shot (<= 4 MiB) and * LSA two-shot ((4 MiB, GIN floor)). * See LICENSE.txt for license information. @@ -61,8 +62,10 @@ int64_t rcclParamGinAllReduceForceEnable(); bool ncclAllReduceGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op); -// True when this AllReduce is a GIN-SDMA candidate that the default size policy -// left for DDA (message below the GIN floor and RCCL_GIN_ALLREDUCE_FORCE_ENABLE is not 1). +// True when this AllReduce is a GIN-SDMA candidate below the GIN floor and +// RCCL_GIN_ALLREDUCE_FORCE_ENABLE is not 1. The floor is 256 MiB on gfx950 and +// 64 MiB on gfx1250, with no lower bound. On gfx1250, sums from symMinR2 +// (8 MiB) up to 64 MiB yield too, so DDA LL/LL128 can claim them. bool ncclAllReduceGinSdmaYieldToDda(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, ncclRedOp_t op); diff --git a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h index 1af758554665..ad20150c95f6 100644 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h @@ -215,6 +215,9 @@ inline bool ginAllReduceSizePolicyEligible(size_t count, size_t typeSize, int nR } // Smaller-than-ginMinBytes GIN candidates yield to DDA unless FORCE_ENABLE=1. +// No lower bound. On gfx1250 (ginMinBytes == 64 MiB) every eligible sum below +// the GinScatter band yields, including (8 MiB, 64 MiB), where DDA LL/LL128 +// can run. Below 8 MiB, symMinR2 already selects DDA. inline bool ginAllReduceYieldToDdaBySize(size_t count, size_t typeSize, bool forceEnable, size_t ginMinBytes = kGinAllReduceGinTwoShotMinBytes) { if (forceEnable) return false; diff --git a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu index f786aabb75d8..0e2b93d00fe6 100755 --- a/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu +++ b/projects/rccl/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu @@ -5,7 +5,9 @@ * Default: gfx950 messages >= 256 MiB take GIN two-shot (LSA reduce-scatter + * GIN all-gather). gfx1250 GinScatter is opt-in * (NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1) for messages in [64 MiB, 4 GiB]. - * Smaller messages fall through to DDA AllReduce. + * Smaller messages yield to DDA. On gfx1250 that yield has no lower bound: + * below symMinR2 (8 MiB) DDA already won, and from there up to 64 MiB DDA + * LL/LL128 can claim the sum. * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables the LSA bands: * <= 4 MiB — LSA one-shot * (4 MiB, GIN floor) — LSA two-shot @@ -93,32 +95,43 @@ static ncclResult_t ginAllReduceScratchInitOnce(ncclComm* comm) { } const size_t allocBytes = static_cast(requested); + // The first GinScatter may be the first collective inside a capture. + // ncclMemAlloc issues cuMemCreate / cuMemMap, which are illegal in global + // capture mode. ncclDevrWindowRegisterInGroup relaxes only around its own + // body, so enter relaxed mode before the allocation. The register call + // nests another exchange and restores relaxed before it returns. Free on + // the failure path stays inside this region for the same reason. + cudaStreamCaptureMode captureMode = cudaStreamCaptureModeRelaxed; + CUDACHECK(cudaThreadExchangeStreamCaptureMode(&captureMode)); + void* scratch = nullptr; ncclWindow_t scratchWinDev = nullptr; ncclResult_t ret = ncclMemAlloc(&scratch, allocBytes); - if (ret != ncclSuccess) return ret; - - ret = ncclDevrWindowRegisterInGroup(comm, scratch, allocBytes, NCCL_WIN_COLL_SYMMETRIC, &scratchWinDev); - if (ret != ncclSuccess) { - (void)ncclMemFree(scratch); - return ret; - } - - ret = ncclDevrFindWindow(comm, scratch, &state->scratchWin); - if (ret != ncclSuccess || scratchWinDev == nullptr || !ginAllReduceScratchWinReady(comm)) { - if (scratchWinDev != nullptr) { - (void)ncclCommWindowDeregister(comm, scratchWinDev); - } else if (state->scratchWin != nullptr && state->scratchWin->vidmem != nullptr) { - (void)ncclCommWindowDeregister(comm, state->scratchWin->vidmem); + if (ret == ncclSuccess) { + ret = ncclDevrWindowRegisterInGroup(comm, scratch, allocBytes, NCCL_WIN_COLL_SYMMETRIC, &scratchWinDev); + if (ret != ncclSuccess) { + (void)ncclMemFree(scratch); + } else { + ret = ncclDevrFindWindow(comm, scratch, &state->scratchWin); + if (ret != ncclSuccess || scratchWinDev == nullptr || !ginAllReduceScratchWinReady(comm)) { + if (scratchWinDev != nullptr) { + (void)ncclCommWindowDeregister(comm, scratchWinDev); + } else if (state->scratchWin != nullptr && state->scratchWin->vidmem != nullptr) { + (void)ncclCommWindowDeregister(comm, state->scratchWin->vidmem); + } + state->scratchWin = nullptr; + (void)ncclMemFree(scratch); + if (ret == ncclSuccess) ret = ncclInternalError; + } else { + state->scratch = scratch; + state->scratchBytes = allocBytes; + INFO(NCCL_INIT, "GIN AllReduce GinScatter scratch: %zu bytes", allocBytes); + } } - state->scratchWin = nullptr; - (void)ncclMemFree(scratch); - return ret != ncclSuccess ? ret : ncclInternalError; } - state->scratch = scratch; - state->scratchBytes = allocBytes; - INFO(NCCL_INIT, "GIN AllReduce GinScatter scratch: %zu bytes", allocBytes); - return ncclSuccess; + + CUDACHECKIGNORE(cudaThreadExchangeStreamCaptureMode(&captureMode)); + return ret; } // Runs on the first eligible AllReduce, which may itself be inside a graph capture. Everything diff --git a/projects/rccl/src/rccl_wrap.cc b/projects/rccl/src/rccl_wrap.cc index b109422c5e85..2f08135728db 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -1667,10 +1667,11 @@ ncclResult_t rcclSelectAllReduce(struct ncclComm* comm, const void* sendbuff, vo // GIN AllReduce is selected first in this function and requires symmetric // windows. By default gfx950 claims messages >= 256 MiB. gfx1250 claims // [64 MiB, 4 GiB] only when NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1, so DDA must - // still be allowed for smaller - // symmetric AllReduces (otherwise they would hit the symmetric kernel - // instead of DDA). FORCE_ENABLE=1 keeps the original - // !symEligible gate because GIN already returned above for those sizes. + // still be allowed for smaller symmetric AllReduces (otherwise they would + // hit the symmetric kernel instead of DDA). The yield has no lower bound. + // Below symMinR2 (8 MiB) DDA already won. From there up to 64 MiB, clearing + // ddaSymEligible lets DDA LL/LL128 claim the sum. FORCE_ENABLE=1 keeps the + // original gate because GIN already returned above for those sizes. // symSuppressedByMin: DDA wins below symMinR2[AR]; do not block it with symkRequested. bool ddaSymEligible = symkRequested && !symSuppressedByMin; #if defined(ENABLE_ROCSHMEM_GIN) diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 29686357604a..15315bca2b58 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -694,8 +694,9 @@ if(BUILD_TESTS) set_tests_properties(rccl-UnitTestsGinSdmaPolicy PROPERTIES LABELS "unit") # GinScatter local reduce on one GPU. Runs ginScatterReduceChunk against a fake - # scratch and a host sum, for staged and full-column layouts. No librccl. - # Skips when no GPU is visible. + # scratch and a host sum, for staged and full-column layouts, including + # multi-rank sums at 64 MiB and 1 GiB. No librccl. Skips when no GPU is visible. + # Registered in the test-runner configs so CI runs it with the other unit tests. add_executable(rccl-UnitTestsGinScatterReduceGpu gin/gin_scatter_reduce_gpu_test.cpp) set_source_files_properties(gin/gin_scatter_reduce_gpu_test.cpp PROPERTIES LANGUAGE CXX) target_include_directories(rccl-UnitTestsGinScatterReduceGpu PRIVATE diff --git a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp index 5db7e27ff307..3a8c7bcb267f 100644 --- a/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp +++ b/projects/rccl/test/gin/gin_all_reduce_policy_test.cpp @@ -126,8 +126,13 @@ TEST(GinAllReducePolicy, Gfx1250ScatterMinBytesFrom64MiB) { } TEST(GinAllReducePolicy, YieldToDdaBySizeGfx1250) { + // No lower bound. Below 8 MiB, symMinR2 already selects DDA. From 8 MiB to + // 64 MiB the same yield lets DDA LL/LL128 take the sum. + EXPECT_TRUE(ginAllReduceYieldToDdaBySize(countForBytes(4 * kMiB), kFloat, false, kGinAllReduceGinScatterMinBytes)); + EXPECT_TRUE(ginAllReduceYieldToDdaBySize(countForBytes(32 * kMiB), kFloat, false, kGinAllReduceGinScatterMinBytes)); EXPECT_TRUE(ginAllReduceYieldToDdaBySize(countForBytes(63 * kMiB), kFloat, false, kGinAllReduceGinScatterMinBytes)); EXPECT_FALSE(ginAllReduceYieldToDdaBySize(countForBytes(64 * kMiB), kFloat, false, kGinAllReduceGinScatterMinBytes)); + EXPECT_FALSE(ginAllReduceYieldToDdaBySize(countForBytes(32 * kMiB), kFloat, true, kGinAllReduceGinScatterMinBytes)); } TEST(GinAllReducePolicy, GinScatterDispatchGfx1250From64MiB) { diff --git a/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp index 37ccf09d8d71..cc02e9d3e900 100644 --- a/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp +++ b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp @@ -7,11 +7,13 @@ // Single-GPU check of ginScatterReduceChunk. Fills a fake incoming scratch // with known per-peer patterns and compares the reduced slice to a host sum. // Covers staged and full-column layouts, including a chunk large enough that -// the 56x512 unroll used to reduce the tail twice. No librccl and no GIN -// communicator. Skips when no GPU is visible. +// the 56x512 unroll used to reduce the tail twice, and multi-rank sums at the +// 64 MiB and 1 GiB GinScatter sizes. No librccl and no GIN communicator. +// Skips when no GPU is visible. #include +#include #include #include #include @@ -112,6 +114,112 @@ void runReduceCase(bool staged, size_t elemStart, size_t nVec, int rank, int nRa HIP_OK(hipFree(dOut)); } +struct DeviceFree { + void* ptr = nullptr; + ~DeviceFree() { + if (ptr != nullptr) hipFree(ptr); + } +}; + +float peerPattern(int peer, size_t elem) { + return static_cast((peer + 1) * 17 + static_cast(elem & 63)); +} + +__global__ void fillRankPattern(float* dst, size_t n, int peer, size_t elemBase) { + const size_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (i >= n) return; + dst[i] = static_cast((peer + 1) * 17 + static_cast((elemBase + i) & 63)); +} + +void launchFill(float* dst, size_t n, int peer, size_t elemBase) { + if (n == 0 || dst == nullptr) return; + constexpr int kThreads = 256; + const unsigned blocks = static_cast((n + static_cast(kThreads) - 1) / kThreads); + fillRankPattern<<>>(dst, n, peer, elemBase); +} + +// One ginScatterReduceChunk at a production size. strideBytes is the incoming +// column (slotBytes when staged, per-rank bytes when not). The chunk itself +// may be shorter than that column. +void runProductionSum(int nRanks, size_t elemStart, size_t chunkBytes, size_t strideBytes, bool staged) { + constexpr int W = static_cast(sizeof(uint4) / sizeof(float)); + ASSERT_EQ(chunkBytes % sizeof(uint4), 0u); + ASSERT_EQ(strideBytes % sizeof(float), 0u); + ASSERT_EQ(elemStart % static_cast(W), 0u); + const size_t chunkElems = chunkBytes / sizeof(float); + const size_t strideElems = strideBytes / sizeof(float); + const size_t elemEnd = elemStart + chunkElems; + ASSERT_GE(strideElems, staged ? chunkElems : elemEnd); + + DeviceFree self; + DeviceFree incoming; + DeviceFree out; + hipError_t err = hipMalloc(&self.ptr, elemEnd * sizeof(float)); + if (err != hipSuccess) { + ADD_FAILURE() << "hipMalloc self: " << hipGetErrorString(err); + (void)hipGetLastError(); + return; + } + err = hipMalloc(&incoming.ptr, static_cast(nRanks) * strideElems * sizeof(float)); + if (err != hipSuccess) { + ADD_FAILURE() << "hipMalloc incoming: " << hipGetErrorString(err); + (void)hipGetLastError(); + return; + } + err = hipMalloc(&out.ptr, elemEnd * sizeof(float)); + if (err != hipSuccess) { + ADD_FAILURE() << "hipMalloc out: " << hipGetErrorString(err); + (void)hipGetLastError(); + return; + } + + float* dSelf = static_cast(self.ptr); + float* dIncoming = static_cast(incoming.ptr); + float* dOut = static_cast(out.ptr); + // Poison the incoming columns. The self column must not be read. + HIP_OK(hipMemset(dIncoming, 0x7f, static_cast(nRanks) * strideElems * sizeof(float))); + // The kernel indexes self by absolute element, so the buffer covers + // [0, elemEnd) and only the chunk range holds this rank's pattern. + const int rank = nRanks > 1 ? 1 : 0; + launchFill(dSelf + elemStart, chunkElems, rank, elemStart); + for (int peer = 0; peer < nRanks; ++peer) { + if (peer == rank) continue; + launchFill(dIncoming + static_cast(peer) * strideElems, strideElems, peer, staged ? elemStart : 0); + } + HIP_OK(hipMemset(dOut, 0, elemEnd * sizeof(float))); + HIP_OK(hipGetLastError()); + + ginScatterReduceChunkKernel<<>>( + dSelf, dIncoming, dOut, rank, nRanks, strideBytes, elemStart, elemEnd, staged); + HIP_OK(hipGetLastError()); + HIP_OK(hipDeviceSynchronize()); + + if (elemStart > 0) { + float head = 1.f; + float before = 1.f; + HIP_OK(hipMemcpy(&head, dOut, sizeof(float), hipMemcpyDeviceToHost)); + HIP_OK(hipMemcpy(&before, dOut + elemStart - 1, sizeof(float), hipMemcpyDeviceToHost)); + EXPECT_EQ(head, 0.f); + EXPECT_EQ(before, 0.f); + } + + constexpr size_t kTile = 1u << 20; + std::vector got(std::min(kTile, chunkElems)); + size_t wrong = 0; + for (size_t base = elemStart; base < elemEnd; base += kTile) { + const size_t n = std::min(kTile, elemEnd - base); + HIP_OK(hipMemcpy(got.data(), dOut + base, n * sizeof(float), hipMemcpyDeviceToHost)); + for (size_t i = 0; i < n; ++i) { + const size_t elem = base + i; + float exp = 0.f; + for (int peer = 0; peer < nRanks; ++peer) exp += peerPattern(peer, elem); + if (got[i] != exp) ++wrong; + } + } + EXPECT_EQ(wrong, 0u) << "staged=" << staged << " elemStart=" << elemStart << " chunkBytes=" << chunkBytes + << " strideBytes=" << strideBytes << " nRanks=" << nRanks; +} + } // namespace TEST(GinScatterReduceGpu, ChunkMatchesHostSum) { @@ -130,3 +238,40 @@ TEST(GinScatterReduceGpu, ChunkMatchesHostSum) { runReduceCase(/*staged=*/true, /*elemStart=*/256, overlapping, /*rank=*/0, /*nRanks=*/4, /*extraStrideElems=*/64); runReduceCase(/*staged=*/false, /*elemStart=*/0, /*nVec=*/100, /*rank=*/1, /*nRanks=*/4, /*extraStrideElems=*/32); } + +TEST(GinScatterReduceGpu, MultiRankSumAt64MiBAnd1GiB) { + if (!gpuAvailable()) GTEST_SKIP() << "no visible GPU"; + size_t freeBytes = 0; + size_t totalBytes = 0; + // Last 1 GiB slot keeps a 128 MiB output plus 256 MiB of incoming columns. + if (hipMemGetInfo(&freeBytes, &totalBytes) != hipSuccess || freeBytes < (640ull << 20)) { + GTEST_SKIP() << "need about 640 MiB free GPU memory for the 64 MiB and 1 GiB sums"; + } + + constexpr int kRanks = 8; + constexpr size_t kScratch = kGinAllReduceTwoShotScratchBytes; + constexpr size_t k64MiB = 64ull * 1024ull * 1024ull; + constexpr size_t k1GiB = 1024ull * 1024ull * 1024ull; + + // 64 MiB fits in one incoming column per rank. 4 and 8 cover the rank loop. + const size_t perRank4 = k64MiB / 4u; + ASSERT_FALSE(ginAllReduceGinScatterStaged(perRank4, 4, kScratch)); + runProductionSum(/*nRanks=*/4, /*elemStart=*/0, perRank4, perRank4, /*staged=*/false); + const size_t perRank8 = k64MiB / static_cast(kRanks); + ASSERT_FALSE(ginAllReduceGinScatterStaged(perRank8, kRanks, kScratch)); + runProductionSum(kRanks, /*elemStart=*/0, perRank8, perRank8, /*staged=*/false); + + // 1 GiB is staged ping-pong. Reduce the first slot and the last slot; the + // last slot is the reuse offset the kernel writes on a later chunk. + const size_t perRank = k1GiB / static_cast(kRanks); + int nChunks = 0; + int nPhases = 0; + size_t uniformBytes = 0; + size_t lastBytes = 0; + ginAllReduceGinScatterChunkPlan(perRank, kRanks, kScratch, &nChunks, &uniformBytes, &lastBytes, &nPhases); + ASSERT_GT(nChunks, 1); + ASSERT_EQ(nPhases, kGinAllReduceGinScatterNumSlots); + runProductionSum(kRanks, /*elemStart=*/0, uniformBytes, uniformBytes, /*staged=*/true); + const size_t lastElem = ginAllReduceGinScatterChunkOff(nChunks - 1, uniformBytes) / sizeof(float); + runProductionSum(kRanks, lastElem, lastBytes, uniformBytes, /*staged=*/true); +} diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json index 4bb851b7c7a4..40a85d492435 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300_mellanox_ib_func.json @@ -2036,6 +2036,21 @@ } ] }, + "unit_tests_gin_scatter_reduce_gpu": { + "is_gtest": true, + "binary": "rccl-UnitTestsGinScatterReduceGpu", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 300, + "tests": [ + { + "name": "GinScatterReduceGpu_AllTests", + "description": "GinScatter local reduce, including 64 MiB and 1 GiB multi-rank sums", + "test_filter": "*" + } + ] + }, "unit_tests_gin_sdma_allgather_policy": { "is_gtest": true, "binary": "rccl-UnitTestsGinSdmaAllGatherPolicy", @@ -3889,6 +3904,12 @@ "config": "unit_tests_gin_sdma_policy", "enabled": true }, + { + "name": "Unit Tests - GinScatter reduce GPU", + "description": "All rccl-UnitTestsGinScatterReduceGpu cases", + "config": "unit_tests_gin_scatter_reduce_gpu", + "enabled": true + }, { "name": "Unit Tests - GIN-SDMA AllGather policy", "description": "All rccl-UnitTestsGinSdmaAllGatherPolicy cases", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json index d18a9522d3a2..d82d4dd393a9 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json @@ -1534,6 +1534,21 @@ } ] }, + "unit_tests_gin_scatter_reduce_gpu": { + "is_gtest": true, + "binary": "rccl-UnitTestsGinScatterReduceGpu", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 300, + "tests": [ + { + "name": "GinScatterReduceGpu_AllTests", + "description": "GinScatter local reduce, including 64 MiB and 1 GiB multi-rank sums", + "test_filter": "*" + } + ] + }, "unit_tests_gin_sdma_allgather_policy": { "is_gtest": true, "binary": "rccl-UnitTestsGinSdmaAllGatherPolicy", @@ -5379,6 +5394,12 @@ "config": "unit_tests_gin_sdma_policy", "enabled": true }, + { + "name": "Unit Tests - GinScatter reduce GPU", + "description": "All rccl-UnitTestsGinScatterReduceGpu cases", + "config": "unit_tests_gin_scatter_reduce_gpu", + "enabled": true + }, { "name": "Unit Tests - GIN-SDMA AllGather policy", "description": "All rccl-UnitTestsGinSdmaAllGatherPolicy cases", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json index 583eb07afda4..443fcf879634 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json @@ -1651,6 +1651,21 @@ } ] }, + "unit_tests_gin_scatter_reduce_gpu": { + "is_gtest": true, + "binary": "rccl-UnitTestsGinScatterReduceGpu", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 300, + "tests": [ + { + "name": "GinScatterReduceGpu_AllTests", + "description": "GinScatter local reduce, including 64 MiB and 1 GiB multi-rank sums", + "test_filter": "*" + } + ] + }, "unit_tests_gin_sdma_allgather_policy": { "is_gtest": true, "binary": "rccl-UnitTestsGinSdmaAllGatherPolicy", @@ -5536,6 +5551,12 @@ "config": "unit_tests_gin_sdma_policy", "enabled": true }, + { + "name": "Unit Tests - GinScatter reduce GPU", + "description": "All rccl-UnitTestsGinScatterReduceGpu cases", + "config": "unit_tests_gin_scatter_reduce_gpu", + "enabled": true + }, { "name": "Unit Tests - GIN-SDMA AllGather policy", "description": "All rccl-UnitTestsGinSdmaAllGatherPolicy cases", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json index 5235b8cb0f72..63334a2f2508 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce_func.json @@ -2104,6 +2104,21 @@ } ] }, + "unit_tests_gin_scatter_reduce_gpu": { + "is_gtest": true, + "binary": "rccl-UnitTestsGinScatterReduceGpu", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 300, + "tests": [ + { + "name": "GinScatterReduceGpu_AllTests", + "description": "GinScatter local reduce, including 64 MiB and 1 GiB multi-rank sums", + "test_filter": "*" + } + ] + }, "unit_tests_gin_sdma_allgather_policy": { "is_gtest": true, "binary": "rccl-UnitTestsGinSdmaAllGatherPolicy", @@ -4060,6 +4075,12 @@ "config": "unit_tests_gin_sdma_policy", "enabled": true }, + { + "name": "Unit Tests - GinScatter reduce GPU", + "description": "All rccl-UnitTestsGinScatterReduceGpu cases", + "config": "unit_tests_gin_scatter_reduce_gpu", + "enabled": true + }, { "name": "Unit Tests - GIN-SDMA AllGather policy", "description": "All rccl-UnitTestsGinSdmaAllGatherPolicy cases", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json index 5d0e8026692d..2ecca429af89 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce.json @@ -2273,6 +2273,21 @@ } ] }, + "unit_tests_gin_scatter_reduce_gpu": { + "is_gtest": true, + "binary": "rccl-UnitTestsGinScatterReduceGpu", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 300, + "tests": [ + { + "name": "GinScatterReduceGpu_AllTests", + "description": "GinScatter local reduce, including 64 MiB and 1 GiB multi-rank sums", + "test_filter": "*" + } + ] + }, "unit_tests_gin_sdma_allgather_policy": { "is_gtest": true, "binary": "rccl-UnitTestsGinSdmaAllGatherPolicy", @@ -7025,6 +7040,12 @@ "config": "unit_tests_gin_sdma_policy", "enabled": true }, + { + "name": "Unit Tests - GinScatter reduce GPU", + "description": "All rccl-UnitTestsGinScatterReduceGpu cases", + "config": "unit_tests_gin_scatter_reduce_gpu", + "enabled": true + }, { "name": "Unit Tests - GIN-SDMA AllGather policy", "description": "All rccl-UnitTestsGinSdmaAllGatherPolicy cases", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json index 14c36b2eefc7..e31cb8d1ae34 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi455_ainic_roce_func.json @@ -2348,6 +2348,21 @@ } ] }, + "unit_tests_gin_scatter_reduce_gpu": { + "is_gtest": true, + "binary": "rccl-UnitTestsGinScatterReduceGpu", + "num_ranks": 1, + "num_nodes": 1, + "num_gpus": 1, + "timeout": 300, + "tests": [ + { + "name": "GinScatterReduceGpu_AllTests", + "description": "GinScatter local reduce, including 64 MiB and 1 GiB multi-rank sums", + "test_filter": "*" + } + ] + }, "unit_tests_gin_sdma_allgather_policy": { "is_gtest": true, "binary": "rccl-UnitTestsGinSdmaAllGatherPolicy", @@ -5617,6 +5632,12 @@ "config": "unit_tests_gin_sdma_policy", "enabled": true }, + { + "name": "Unit Tests - GinScatter reduce GPU", + "description": "All rccl-UnitTestsGinScatterReduceGpu cases", + "config": "unit_tests_gin_scatter_reduce_gpu", + "enabled": true + }, { "name": "Unit Tests - GIN-SDMA AllGather policy", "description": "All rccl-UnitTestsGinSdmaAllGatherPolicy cases", diff --git a/runtimes/ddi/rocddi/src/topology/snapshot.rs b/runtimes/ddi/rocddi/src/topology/snapshot.rs deleted file mode 100644 index 4ae394281cc9..000000000000 --- a/runtimes/ddi/rocddi/src/topology/snapshot.rs +++ /dev/null @@ -1,682 +0,0 @@ -// Copyright (c) 2026 Advanced Micro Devices, Inc. -// SPDX-License-Identifier: MIT - -//! One passive, published view of independently reported topology facts. -//! -//! Node IDs are indices into this snapshot, never native IDs. A KFD node -//! ordinal and a Linux NUMA ID are separate namespaces even when equal. -//! Routes are directed reports by one source, not a physical switch graph or -//! permission to map memory. KFD sources have a checked native generation; -//! the host OS source is best effort and has no atomic generation with KFD. - -use super::host::HostTopologySnapshot; -use super::{CacheInfo, Endpoint, MemoryLinkInfo}; -use crate::driver::DriverInstanceId; -use crate::memory::storage::TryPush; -use crate::{Error, ErrorKind}; - -/// Index valid only in the snapshot that returned it. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct NodeId(pub usize); - -/// Index valid only in the snapshot that returned it. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct SourceId(pub usize); - -/// Source of a group of independently observed topology records. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub enum SourceKind { - /// The Linux host OS CPU and NUMA reader. - HostOs, - /// One installed Linux KFD driver instance. - LinuxKfd, - /// Another installed activation driver. - Driver, -} - -/// Freshness and completeness of one contributor's records. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct TopologySource { - /// Contributor family; source IDs distinguish instances of one family. - pub kind: SourceKind, - /// Native generation checked around this contributor's scan, if any. - pub generation: Option, - /// Whether this contributor read all of its advertised records. - pub complete: bool, - /// Activation owner, absent for the passive host OS contributor. - pub(crate) driver_instance: Option, -} - -/// Native node identities remain qualified by their source and type. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub enum NativeNodeId { - /// Host-wide CPU/cache observation. - HostSystem, - /// OS NUMA node number. - LinuxNuma(u32), - /// KFD topology node ordinal. - LinuxKfd(u32), - /// Opaque ordinal reported by another activation driver. - Driver(u64), -} - -/// Classification of a source node, independent of activated devices. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub enum NodeKind { - /// Host-wide OS observation. - HostSystem, - /// OS NUMA domain. - HostNuma, - /// KFD node with no GPU ID. - KfdHost, - /// KFD node reporting a GPU ID. - KfdGpu, - /// Node reported by another activation driver. - Other, -} - -/// One source-qualified topology node. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct TopologyNode { - /// Snapshot-scoped node index. - pub id: NodeId, - /// Source that reported this node. - pub source: SourceId, - /// Source-qualified native node name. - pub native_id: NativeNodeId, - /// Source node classification. - pub kind: NodeKind, -} - -/// An endpoint can attach to multiple nodes; several endpoints may attach to -/// the same node. This record does not own an activated device. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub enum AttachmentKind { - /// Activation driver node that reported the endpoint. - Native, - /// Best-effort host NUMA locality reported by the OS. - HostLocality, -} - -/// One typed endpoint-to-node association. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct EndpointAttachment { - /// Passive endpoint ID; it does not name a device handle. - pub endpoint_id: [u8; 16], - /// Node on which this endpoint was reported. - pub node: NodeId, - /// Meaning of this node association. - pub kind: AttachmentKind, -} - -/// Directed native route reported by one contributor. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct ReportedRoute { - /// Origin node of this directed report. - pub source: NodeId, - /// Destination node of this directed report. - pub target: NodeId, - /// Native latency, bandwidth, and link-type facts. - pub info: MemoryLinkInfo, -} - -/// Cache associated with a source node, not an activated device. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct TopologyCache { - /// Node at which the source reported the cache. - pub node: NodeId, - /// Passive cache size, level, line size, and kind. - pub info: CacheInfo, -} - -/// Passive placement class for a reported memory capacity. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub enum MemoryLocationKind { - /// Host system memory reported by the OS. - HostSystem, - /// Separate local memory reported by an activation driver. - DeviceLocal, -} - -/// Memory capacity attached to a node, without access authorization. -#[derive(Clone, Copy, Debug, Eq, PartialEq)] -pub struct MemoryLocation { - /// Source node at which the memory was reported. - pub node: NodeId, - /// Host or device-local placement classification. - pub kind: MemoryLocationKind, - /// Total bytes, when the contributor reported a capacity. - pub capacity_bytes: Option, - /// Host-visible subset for device-local memory, when reported. - pub host_visible_bytes: Option, -} - -/// Complete candidate published only after all driver contributions validate. -/// Host facts can be incomplete; their per-source `complete` flag records that. -#[derive(Clone, Debug, Eq, PartialEq)] -pub struct TopologySnapshot { - /// Best-effort OS host observations used for CPU frontend policy. - pub host: HostTopologySnapshot, - /// Passive activation endpoints discovered by installed drivers. - pub endpoints: Vec, - /// Freshness and completeness of each independent contributor. - pub sources: Vec, - /// Source-qualified nodes, including nodes without an endpoint. - pub nodes: Vec, - /// Reported endpoint-to-node associations. - pub attachments: Vec, - /// Directed native routes; no access permission is implied. - pub routes: Vec, - /// Caches attached to source-qualified nodes. - pub caches: Vec, - /// Memory locations attached to source-qualified nodes. - pub memory_locations: Vec, -} - -impl TopologySnapshot { - pub(crate) fn new(host: HostTopologySnapshot) -> Result { - let mut result = Self { - host, - endpoints: Vec::new(), - sources: Vec::new(), - nodes: Vec::new(), - attachments: Vec::new(), - routes: Vec::new(), - caches: Vec::new(), - memory_locations: Vec::new(), - }; - let complete = result.host.domains_complete - && result.host.caches.complete - && result - .host - .domains - .iter() - .all(|domain| domain.caches.complete); - let source = result.add_source(SourceKind::HostOs, None, complete, None)?; - let system = result.add_node(source, NativeNodeId::HostSystem, NodeKind::HostSystem)?; - result.memory_locations.try_push(MemoryLocation { - node: system, - kind: MemoryLocationKind::HostSystem, - capacity_bytes: result.host.memory_bytes.map(|bytes| bytes as u64), - host_visible_bytes: None, - })?; - for cache in &result.host.caches.records { - result.caches.try_push(TopologyCache { - node: system, - info: cache.info, - })?; - } - for domain in &result.host.domains { - if result.nodes.iter().any(|node| { - node.source == source && node.native_id == NativeNodeId::LinuxNuma(domain.numa_node) - }) { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "host source reported a duplicate NUMA node", - }); - } - let node = NodeId(result.nodes.len()); - result.nodes.try_push(TopologyNode { - id: node, - source, - native_id: NativeNodeId::LinuxNuma(domain.numa_node), - kind: NodeKind::HostNuma, - })?; - for cache in &domain.caches.records { - result.caches.try_push(TopologyCache { - node, - info: cache.info, - })?; - } - } - Ok(result) - } - - pub(crate) fn add_source( - &mut self, - kind: SourceKind, - generation: Option, - complete: bool, - driver_instance: Option, - ) -> Result { - if (kind == SourceKind::HostOs) != driver_instance.is_none() { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "topology source has an invalid activation owner", - }); - } - let id = SourceId(self.sources.len()); - self.sources.try_push(TopologySource { - kind, - generation, - complete, - driver_instance, - })?; - Ok(id) - } - - pub(crate) fn add_node( - &mut self, - source: SourceId, - native_id: NativeNodeId, - kind: NodeKind, - ) -> Result { - let valid_type = self.sources.get(source.0).is_some_and(|source| { - matches!( - (source.kind, native_id, kind), - ( - SourceKind::HostOs, - NativeNodeId::HostSystem, - NodeKind::HostSystem - ) | ( - SourceKind::HostOs, - NativeNodeId::LinuxNuma(_), - NodeKind::HostNuma - ) | ( - SourceKind::LinuxKfd, - NativeNodeId::LinuxKfd(_), - NodeKind::KfdHost | NodeKind::KfdGpu - ) | (SourceKind::Driver, NativeNodeId::Driver(_), NodeKind::Other) - ) - }); - if !valid_type - || self - .nodes - .iter() - .any(|node| node.source == source && node.native_id == native_id) - { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "topology source reported a duplicate or invalid node", - }); - } - let id = NodeId(self.nodes.len()); - self.nodes.try_push(TopologyNode { - id, - source, - native_id, - kind, - })?; - Ok(id) - } - - pub(crate) fn add_endpoint(&mut self, endpoint: Endpoint, node: NodeId) -> Result<(), Error> { - if self.nodes.get(node.0).is_none() - || self.sources[self.nodes[node.0].source.0].driver_instance - != Some(endpoint.driver_instance) - || self.endpoints.iter().any(|record| record.id == endpoint.id) - { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "topology endpoint has a duplicate ID or invalid attachment", - }); - } - self.attachments.try_push(EndpointAttachment { - endpoint_id: endpoint.id, - node, - kind: AttachmentKind::Native, - })?; - self.endpoints.try_push(endpoint)?; - Ok(()) - } - - pub(crate) fn attach_endpoint( - &mut self, - endpoint_id: [u8; 16], - driver_instance: DriverInstanceId, - node: NodeId, - kind: AttachmentKind, - ) -> Result<(), Error> { - if self.nodes.get(node.0).is_none() - || !self - .endpoints - .iter() - .any(|record| record.id == endpoint_id && record.driver_instance == driver_instance) - || self.attachments.iter().any(|attachment| { - attachment.endpoint_id == endpoint_id - && attachment.node == node - && attachment.kind == kind - }) - { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "topology attachment is duplicate or names an unknown record", - }); - } - self.attachments.try_push(EndpointAttachment { - endpoint_id, - node, - kind, - })?; - Ok(()) - } - - pub(crate) fn add_route( - &mut self, - source: NodeId, - target: NodeId, - info: MemoryLinkInfo, - ) -> Result<(), Error> { - let from = self.nodes.get(source.0); - let to = self.nodes.get(target.0); - if !matches!((from, to), (Some(a), Some(b)) if a.source == b.source) { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "reported route crosses topology sources or has an invalid node", - }); - } - if let Some(route) = self - .routes - .iter_mut() - .find(|route| route.source == source && route.target == target) - { - route.info = info; - return Ok(()); - } - self.routes.try_push(ReportedRoute { - source, - target, - info, - })?; - Ok(()) - } - - pub(crate) fn add_memory_location(&mut self, location: MemoryLocation) -> Result<(), Error> { - if self.nodes.get(location.node.0).is_none() - || location - .host_visible_bytes - .zip(location.capacity_bytes) - .is_some_and(|(visible, total)| visible > total) - { - return Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "topology memory location is inconsistent or has no node", - }); - } - if let Some(existing) = self - .memory_locations - .iter() - .find(|existing| existing.node == location.node && existing.kind == location.kind) - { - return if *existing == location { - Ok(()) - } else { - Err(Error::Operation { - kind: ErrorKind::InvalidData, - detail: "topology source reported conflicting memory capacities", - }) - }; - } - self.memory_locations.try_push(location)?; - Ok(()) - } - - /// Nodes to which the endpoint is attached in this snapshot. - pub fn nodes_for_endpoint(&self, endpoint: &Endpoint) -> impl Iterator { - let owned = self.endpoints.iter().any(|candidate| { - candidate.id == endpoint.id && candidate.driver_instance == endpoint.driver_instance - }); - self.attachments - .iter() - .filter(move |attachment| owned && attachment.endpoint_id == endpoint.id) - .filter_map(|attachment| self.nodes.get(attachment.node.0)) - } - - fn native_nodes_for_endpoint( - &self, - endpoint: &Endpoint, - ) -> impl Iterator { - self.nodes_for_endpoint(endpoint).filter(|node| { - self.attachments.iter().any(|attachment| { - attachment.endpoint_id == endpoint.id - && attachment.node == node.id - && attachment.kind == AttachmentKind::Native - }) - }) - } - - /// Cache records reported by the endpoint's native activation nodes. - /// Host-locality attachments do not make host caches into GPU caches. - pub fn caches_for_endpoint(&self, endpoint: &Endpoint) -> impl Iterator { - self.caches.iter().filter_map(move |cache| { - self.native_nodes_for_endpoint(endpoint) - .any(|node| node.id == cache.node) - .then_some(&cache.info) - }) - } - - /// Reported directed route between two endpoint attachments. - #[must_use] - pub fn route_between(&self, from: &Endpoint, to: &Endpoint) -> Option { - self.routes.iter().rev().find_map(|route| { - let source = self.nodes.get(route.source.0)?; - let target = self.nodes.get(route.target.0)?; - (source.source == target.source - && self - .native_nodes_for_endpoint(from) - .any(|node| node.id == source.id) - && self - .native_nodes_for_endpoint(to) - .any(|node| node.id == target.id)) - .then_some(route.info) - }) - } - - /// Reported directed route from an endpoint to a KFD host node. - #[must_use] - pub fn route_to_host(&self, from: &Endpoint) -> Option { - self.routes.iter().rev().find_map(|route| { - let source = self.nodes.get(route.source.0)?; - let target = self.nodes.get(route.target.0)?; - (source.source == target.source - && target.kind == NodeKind::KfdHost - && self - .native_nodes_for_endpoint(from) - .any(|node| node.id == source.id)) - .then_some(route.info) - }) - } - - /// Reported directed route from a KFD host node to an endpoint. - #[must_use] - pub fn route_from_host(&self, to: &Endpoint) -> Option { - self.routes.iter().rev().find_map(|route| { - let source = self.nodes.get(route.source.0)?; - let target = self.nodes.get(route.target.0)?; - (source.source == target.source - && source.kind == NodeKind::KfdHost - && self - .native_nodes_for_endpoint(to) - .any(|node| node.id == target.id)) - .then_some(route.info) - }) - } -} - -#[cfg(test)] -#[allow(clippy::unwrap_used)] -mod tests { - use super::*; - use crate::driver::{EndpointSelector, new_driver_instance}; - use crate::memory::DeviceAccess; - use crate::topology::host::{HostCaches, HostDomain}; - use crate::topology::{CacheKind, EndpointKind}; - - fn cache(level: u32) -> CacheInfo { - CacheInfo::new(level, 4096, Some(64), CacheKind::Unified) - } - - fn host() -> HostTopologySnapshot { - HostTopologySnapshot { - memory_bytes: None, - cpu: None, - online_cpus: None, - caches: HostCaches { - records: Vec::new(), - complete: false, - }, - domains: vec![HostDomain { - numa_node: 7, - cpus: None, - caches: HostCaches { - records: Vec::new(), - complete: true, - }, - }], - domains_complete: true, - } - } - - fn endpoint(id: u8) -> Endpoint { - Endpoint { - id: [id; 16], - name: [0; 128], - product_name: None, - pci: None, - kind: EndpointKind::Npu, - local_memory_bytes: 0, - host_visible_local_memory_bytes: 0, - host_local_cacheability: None, - allocation_granularity: 4096, - address_bit_count: 48, - minimum_address: 0, - maximum_address: u64::MAX, - supported_permissions: DeviceAccess::NONE, - driver_instance: new_driver_instance(), - selector: EndpointSelector::Opaque(u64::from(id)), - } - } - - fn link() -> MemoryLinkInfo { - MemoryLinkInfo { - minimum_latency: 1, - maximum_latency: 2, - minimum_bandwidth: 3, - maximum_bandwidth: 4, - atomic_support_32bit: false, - atomic_support_64bit: false, - coherent_support: false, - kind: super::super::MemoryLinkType::Pcie, - numa_distance: 20, - } - } - - #[test] - fn source_qualified_nodes_allow_multiple_endpoints_and_typed_attachments() { - let mut snapshot = TopologySnapshot::new(host()).unwrap(); - let first = endpoint(1); - let mut second = endpoint(2); - second.driver_instance = first.driver_instance; - let source = snapshot - .add_source( - SourceKind::LinuxKfd, - Some(91), - true, - Some(first.driver_instance), - ) - .unwrap(); - let gpu = snapshot - .add_node(source, NativeNodeId::LinuxKfd(7), NodeKind::KfdGpu) - .unwrap(); - let host_numa = snapshot - .nodes - .iter() - .find(|node| node.native_id == NativeNodeId::LinuxNuma(7)) - .unwrap() - .id; - assert_ne!(gpu, host_numa); - assert_eq!(snapshot.nodes[gpu.0].source, source); - assert_eq!(snapshot.nodes[host_numa.0].source, SourceId(0)); - assert_eq!(snapshot.sources[0].generation, None); - assert!(!snapshot.sources[0].complete); - assert_eq!(snapshot.sources[source.0].generation, Some(91)); - snapshot.add_endpoint(first.clone(), gpu).unwrap(); - snapshot.add_endpoint(second.clone(), gpu).unwrap(); - snapshot - .attach_endpoint( - first.id, - first.driver_instance, - host_numa, - AttachmentKind::HostLocality, - ) - .unwrap(); - snapshot.caches.push(TopologyCache { - node: gpu, - info: cache(2), - }); - snapshot.caches.push(TopologyCache { - node: host_numa, - info: cache(3), - }); - let location = MemoryLocation { - node: gpu, - kind: MemoryLocationKind::DeviceLocal, - capacity_bytes: Some(8192), - host_visible_bytes: Some(4096), - }; - snapshot.add_memory_location(location).unwrap(); - snapshot.add_memory_location(location).unwrap(); - assert_eq!(snapshot.memory_locations.len(), 2); - assert_eq!(snapshot.memory_locations[1], location); - assert_eq!( - snapshot - .add_memory_location(MemoryLocation { - capacity_bytes: Some(16384), - ..location - }) - .unwrap_err() - .kind(), - ErrorKind::InvalidData - ); - assert_eq!(snapshot.nodes_for_endpoint(&first).count(), 2); - assert_eq!(snapshot.nodes_for_endpoint(&second).count(), 1); - assert_eq!(snapshot.caches_for_endpoint(&first).count(), 1); - assert_eq!( - snapshot.caches_for_endpoint(&first).next().unwrap().level, - 2 - ); - let mut foreign = first.clone(); - foreign.driver_instance = new_driver_instance(); - assert_eq!(snapshot.nodes_for_endpoint(&foreign).count(), 0); - } - - #[test] - fn routes_are_directed_reports_with_no_cross_source_inference() { - let mut snapshot = TopologySnapshot::new(host()).unwrap(); - let first = endpoint(1); - let mut second = endpoint(2); - second.driver_instance = first.driver_instance; - let source = snapshot - .add_source( - SourceKind::LinuxKfd, - Some(8), - true, - Some(first.driver_instance), - ) - .unwrap(); - let from = snapshot - .add_node(source, NativeNodeId::LinuxKfd(1), NodeKind::KfdGpu) - .unwrap(); - let to = snapshot - .add_node(source, NativeNodeId::LinuxKfd(2), NodeKind::KfdGpu) - .unwrap(); - let host_node = snapshot - .add_node(source, NativeNodeId::LinuxKfd(0), NodeKind::KfdHost) - .unwrap(); - snapshot.add_endpoint(first.clone(), from).unwrap(); - snapshot.add_endpoint(second.clone(), to).unwrap(); - snapshot.add_route(from, to, link()).unwrap(); - snapshot.add_route(host_node, from, link()).unwrap(); - assert_eq!(snapshot.route_between(&first, &second), Some(link())); - assert_eq!(snapshot.route_between(&second, &first), None); - assert_eq!(snapshot.route_from_host(&first), Some(link())); - assert_eq!(snapshot.route_to_host(&first), None); - assert_eq!( - snapshot - .add_route(from, NodeId(1), link()) - .unwrap_err() - .kind(), - ErrorKind::InvalidData - ); - } -} From d75c3a90ba7cd9c468fa30d9756e21263ff8120b Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Thu, 29 Oct 2026 15:50:02 -0500 Subject: [PATCH 6/7] resolve merge conflict --- .../rocddi/src/driver/linux_kfd/tests/vmem.rs | 33 ------------------- 1 file changed, 33 deletions(-) diff --git a/runtimes/ddi/rocddi/src/driver/linux_kfd/tests/vmem.rs b/runtimes/ddi/rocddi/src/driver/linux_kfd/tests/vmem.rs index 8afc13f3ce8a..f1b5549efb92 100644 --- a/runtimes/ddi/rocddi/src/driver/linux_kfd/tests/vmem.rs +++ b/runtimes/ddi/rocddi/src/driver/linux_kfd/tests/vmem.rs @@ -171,39 +171,6 @@ fn extended_coherence_rejects_uncached_request_from_either_source() { ); } -#[test] -fn extended_coherence_rejects_uncached_request_from_either_source() { - let (vm, _address, _physical) = fixture(); - let kind = MemoryKind::DeviceLocal { - host_visible: false, - coherent: true, - extended_coherent: true, - uncached: false, - contiguous: false, - }; - assert_eq!( - KfdVirtualMemory::create(vm.clone(), kind, page(), false, true) - .err() - .unwrap() - .kind(), - ErrorKind::InvalidArgument - ); - let kind_uncached = MemoryKind::DeviceLocal { - host_visible: false, - coherent: true, - extended_coherent: true, - uncached: true, - contiguous: false, - }; - assert_eq!( - KfdVirtualMemory::create(vm, kind_uncached, page(), false, false) - .err() - .unwrap() - .kind(), - ErrorKind::InvalidArgument - ); -} - #[test] fn ambiguous_map_retains_address_backing_gem_and_pending_point() { let (vm, mut address, mut physical) = fixture(); From f0c1364a833d485c5ebf48cf9565bff0ccb53b68 Mon Sep 17 00:00:00 2001 From: Nusrat Islam Date: Thu, 29 Oct 2026 16:40:04 -0500 Subject: [PATCH 7/7] update eligibility test --- .../test/GinAllReduceEligibilityTests.cpp | 118 ++++++++++++++++++ 1 file changed, 118 insertions(+) diff --git a/projects/rccl/test/GinAllReduceEligibilityTests.cpp b/projects/rccl/test/GinAllReduceEligibilityTests.cpp index cde7422acda9..1a7a8daec167 100644 --- a/projects/rccl/test/GinAllReduceEligibilityTests.cpp +++ b/projects/rccl/test/GinAllReduceEligibilityTests.cpp @@ -5,10 +5,17 @@ ************************************************************************/ #include "common/GinAllReduceTestHelpers.hpp" +#include "common/ProcessIsolatedTestRunner.hpp" #include "algorithms/gin/gin_all_reduce.h" +#include "nccl_device/net_device.h" + #include "gtest/gtest.h" +#include +#include +#include + namespace RcclUnitTesting { @@ -71,6 +78,117 @@ TEST_F(GinAllReduceEligibilityTest, NoSymmetricSupportRejected) EXPECT_FALSE(ncclAllReduceGinSdmaYieldToDda(mockComm_.get(), sendbuff_, recvbuff_, kSmallCount, ncclFloat32, ncclSum)); } +// Stand-in that passes every gate after the gfx1250 opt-in, so deleting that +// check flips the disabled case from false to true. symkState.initialized and +// devrState.bigSize skip ncclSymkInitOnce / ncclDevrInitOnce. The sorted-window +// record matches ncclDevrWindowSorted in src/dev_runtime.cc (private to that file). +struct Gfx1250EligibleStandIn +{ + ncclComm comm{}; + ncclSharedResources shared{}; + ncclDevrWindow win{}; + struct SortedWin + { + uintptr_t userAddr; + size_t size; + ncclDevrWindow* win; + } sorted{}; + char arch_[16]{}; + + static constexpr uintptr_t kBase = 0x1000; + static constexpr size_t kSpan = 0x10000; + + Gfx1250EligibleStandIn() + { + static_assert(offsetof(SortedWin, userAddr) == 0, "window sort key must lead"); + static_assert(offsetof(SortedWin, size) == sizeof(uintptr_t), "window sort size offset"); + static_assert(offsetof(SortedWin, win) == sizeof(uintptr_t) + sizeof(size_t), "window sort pointer offset"); + + std::memcpy(arch_, "gfx1250", 7); + comm.archName = arch_; + comm.nNodes = 1; + comm.nRanks = 8; + comm.rank = 0; + comm.symmetricSupport = true; + comm.isAllDirectNvlink = true; + comm.globalGinSupport = NCCL_GIN_CONNECTION_FULL; + comm.symkState.initialized = true; + comm.devrState.bigSize = 1; + comm.devrState.lsaSize = comm.nRanks; + comm.devrState.lsaSelf = 0; + comm.sharedRes = &shared; + + shared.ginState.finalized = false; + shared.ginState.numActiveBackends = 1; + shared.ginState.backends[0].ginType = static_cast(NCCL_NET_DEVICE_GIN_ANVIL_SDMA); + + win.winFlags = NCCL_WIN_COLL_SYMMETRIC; + win.size = kSpan; + sorted.userAddr = kBase; + sorted.size = kSpan; + sorted.win = &win; + comm.devrState.winSorted = + reinterpret_cast(&sorted); + comm.devrState.winSortedCount = 1; + } + + ncclComm* get() { return &comm; } + void* send() const { return reinterpret_cast(kBase); } + void* recv() const { return reinterpret_cast(kBase + (kSpan / 2)); } +}; + +// NCCL_GIN_ALLREDUCE_GFX1250_ENABLE is cached on first read, so each setting +// runs in its own process. Default off rejects gfx1250 even in the GinScatter +// band. On: [64 MiB, 4 GiB] is eligible, below 64 MiB yields to DDA, above +// 4 GiB is left to the other selectors. +TEST_F(GinAllReduceEligibilityTest, Gfx1250OptIn) +{ + auto countForBytes = [](size_t bytes) { return bytes / sizeof(float); }; + const size_t scatterCount = countForBytes(kGinAllReduceGinScatterMinBytes); + const size_t belowCount = countForBytes(kGinAllReduceLsaOneShotMaxBytes); + const size_t aboveCount = countForBytes(kGinAllReduceGinScatterMaxBytes + 256); + + auto checkDisabled = [&]() { + Gfx1250EligibleStandIn standIn; + EXPECT_FALSE(ncclAllReduceGinSdmaEligible(standIn.get(), standIn.send(), standIn.recv(), scatterCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaYieldToDda(standIn.get(), standIn.send(), standIn.recv(), scatterCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaEligible(standIn.get(), standIn.send(), standIn.recv(), belowCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaYieldToDda(standIn.get(), standIn.send(), standIn.recv(), belowCount, + ncclFloat32, ncclSum)); + }; + auto checkEnabled = [&]() { + Gfx1250EligibleStandIn standIn; + EXPECT_TRUE(ncclAllReduceGinSdmaEligible(standIn.get(), standIn.send(), standIn.recv(), scatterCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaYieldToDda(standIn.get(), standIn.send(), standIn.recv(), scatterCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaEligible(standIn.get(), standIn.send(), standIn.recv(), belowCount, + ncclFloat32, ncclSum)); + EXPECT_TRUE(ncclAllReduceGinSdmaYieldToDda(standIn.get(), standIn.send(), standIn.recv(), belowCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaEligible(standIn.get(), standIn.send(), standIn.recv(), aboveCount, + ncclFloat32, ncclSum)); + EXPECT_FALSE(ncclAllReduceGinSdmaYieldToDda(standIn.get(), standIn.send(), standIn.recv(), aboveCount, + ncclFloat32, ncclSum)); + }; + + auto cleared = [](ProcessIsolatedTestRunner::TestConfig cfg) { + cfg.clearVariable("NCCL_GIN_ALLREDUCE_ENABLE") + .clearVariable("NCCL_GIN_ALLREDUCE_SCRATCH_BYTES") + .clearVariable("RCCL_GIN_ALLREDUCE_FORCE_ENABLE"); + return cfg; + }; + + RUN_ISOLATED_TESTS( + cleared(ProcessIsolatedTestRunner::TestConfig("Gfx1250DisabledByDefault", checkDisabled)) + .clearVariable("NCCL_GIN_ALLREDUCE_GFX1250_ENABLE"), + cleared(ProcessIsolatedTestRunner::TestConfig("Gfx1250EnabledBands", checkEnabled)) + .setVariable("NCCL_GIN_ALLREDUCE_GFX1250_ENABLE", "1")); +} + #endif // ENABLE_ROCSHMEM_GIN #if !defined(ENABLE_ROCSHMEM_GIN)