diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 760d50d3ae2b..d4bc4e422dd5 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -44,6 +44,8 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `RCCL_FABRIC_TELEMETRY_INTERVAL_MS` (default `1000`): sampling period for `RCCL_FABRIC_TELEMETRY_ENABLE`, clamped to a 100 ms floor and a one-hour ceiling. A non-positive value disables telemetry entirely. * Hierarchical symmetric `ncclAllGather` for multi-node using GIN and LSA (`AllGather_RailRing_LsaST`): a GIN ring over the rails team moves data between nodes and LSA stores broadcast it to the other GPUs in each node. * 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. * `RCCL_TUNER_EMBEDDED_CONFIG` (default `1`): the built-in CSV tuner configurations in `tuner/*.csv` are now compiled into `librccl.so` and used as the stock tuning defaults, so a cluster no longer has to stage the CSV on every node or on NFS. Set `0` to ignore the embedded configurations. `NCCL_TUNER_CONFIG_FILE` and the on-disk tuner directories still take precedence, and `NCCL_TUNER_PLUGIN=none` still disables the tuner entirely. * `RCCL_TUNING_PARTITION_MODE` (`auto`, `SPX` or `DPX`; default `auto`): selects the gfx1250 tuning table for the GPU's compute-partition mode. `auto` uses the DPX table when the GPU reports DPX, and the SPX table otherwise. A communicator uses the DPX table only when every rank selects it. The DPX table is currently a copy of the SPX table, so tuning is unchanged. * `RCCL_CE_AR_2SHOT_MAX_BYTES` (default `-1`): overrides the 2-shot AllReduce size cap from the arch table. Set to a positive value to override `ceNonRegMax[AR]`. Renamed from `RCCL_CE_AR_MAX_MSG_BYTES`. @@ -107,6 +109,60 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `NCCL_SYM_TMA_ENABLE` now defaults to `1`, matching NCCL 2.31.2, so on gfx1250 (MI450) the Tensor Data Mover symmetric kernels (`AllGather_TmaST`, `AllReduce_RSxTmaLD_AGxTmaST`, `ReduceScatter_TmaLD`) are offered to the symmetric cost model whenever an application uses symmetric registered memory. The cost model selects them only where their deep loop engages, and launches them at the widest allowed block count. They require 16-byte aligned buffers and a build with gfx1250 in `GPU_TARGETS` and a ROCm SDK that ships the TDM descriptor header. For in-place AllGather and ReduceScatter the per-rank slice must also be a multiple of 16 bytes, so that every rank selects the same kernel. Other architectures are unaffected. Set `NCCL_SYM_TMA_ENABLE=0` to restore the previous behavior. * rccl-tests `-A 1` (and `rcclGetCollImplInfo` / `rcclSymKGetInfo`) now report a symmetric collective that selects a Tensor Data Mover kernel as `SymTDM` instead of `SYM`. The symmetric reporting query also accounts for buffer alignment, so it no longer omits the TDM kernels that the live dispatch would pick. * The AMD Pensando 'ionic' RDMA provider for the rocSHMEM GDA GIN backend (`NCCL_GIN_TYPE=6`) now properly enforces the required support for creation of 'collapsed' single-entry completion queues. If network driver and firmware support for this feature is unavailable, an initialization failure will occur rather than a silent fallback to plain completion queues. +### Removed +* `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it. +* `RCCL_DDA_LL128_RS_THREADS` removed (was default `1024`): reduce-scatter LL128 now launches a fixed 512-thread block, one warp per slice, so the thread count is no longer a tuning knob. + +### Resolved issues +* Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU. +* Fixed profiler overhead when no profiler plugin is loaded or only kernel-channel events are enabled. +* Fixed profiler API events reporting rank 0 instead of the originating communicator rank. +* Fixed P2P IPC registration reuse producing out-of-bounds remote addresses when a registered allocation spans multiple cuMem segments. +* Fixed non-thread-safe token parsing that could corrupt configuration parsed concurrently from multiple threads. +* Fixed tuner plugins receiving uninitialized cost-model constants. +* Fixed GIN proxy descriptor shared-memory sizing and alignment. +* Fixed virtual address space exhaustion when symmetric windows backed by the same physical allocation are registered repeatedly. +* Fixed `NCCL_P2P_DISABLE` and `NCCL_P2P_LEVEL` being ignored on Intel, ARM and Zhaoxin hosts, where the CPU-based default P2P level overwrote the user setting. On such hosts `NCCL_P2P_DISABLE=1` now falls back to SHM as documented, and `NCCL_P2P_LEVEL` is honored. +* Fixed Copy Engine `ncclAlltoAllv` hanging on intra-node communicators of more than 8 ranks with uneven send sizes. Each rank decided from its own send sizes whether to add intra-batch barriers, so ranks could issue different barrier counts. The decision now uses the full send-size matrix, which every rank shares. +* Fixed LL-protocol latency on gfx950 when RCCL is built against a ROCm older than 7.0.2 without the 128-bit global load/store builtins (for example ROCm 7.0.1, or `-DDWORDX4_INTRINSICS=OFF`). Every store to the user output buffer was followed by a release fence and an L2 writeback, which made 4 KB to 64 KB `ncclBroadcast` about 30% slower on MI355X. Builds against ROCm 7.0.2 or later were not affected and are unchanged. +* Fixed a 15 to 30% `ncclScatter` bandwidth loss from 16 MB upward on multi-node AINIC jobs, caused by the move to two network channels per peer (for example on 4 nodes of 8 GPUs). Scatter now uses and connects the most channels per peer whose flows still fit the P2P channel pool. Point-to-point operations and the other collectives keep their channel count. +* Fixed intermittent data corruption in multi-node collectives, seen as out-of-bounds values in 8-node AllToAll, on hosts that use GPU Direct RDMA through peermem. Since RCCL 2.30.4 the GDR flush scratchpad was registered only through dma-buf, so on peermem hosts the flush degraded to a single RDMA read of the receive buffer, which does not order earlier relaxed-ordering writes. The scratchpad is now also registered through peermem, restoring the strictly ordered RDMA write that precedes the flush read. +* Fixed `ncclCommAbort` and `ncclCommRevoke` hanging when a DDA collective (IPC on gfx942/gfx950, fabric LL/LL128/VMM on gfx1250) was waiting on a peer that never joined. DDA kernels now observe the communicator abort flag while waiting. A parent communicator shrunk with `NCCL_SHRINK_ABORT` no longer uses DDA for later collectives. + +### Known issues +* The FP8 ReduceSum and ReduceCopy device APIs are not exported in the LLVM bitcode library. +* On the Anvil SDMA (`NCCL_GIN_TYPE=7`) and rocSHMEM GDA (`NCCL_GIN_TYPE=6`) GIN backends, `ncclGinFlushTimeout` ignores the timeout and blocks until the flush completes; `ncclGinWaitTimeout` does the same on Anvil SDMA. rocSHMEM GDA does not support `ncclGinFlushAsync` or waiting on a GIN request, and those calls trap on the device. The signal and counter wait timeouts work on all backends. + +## RCCL 2.31.2 for ROCm 10.2.0 (Unreleased) + +### Added +* `RCCL_CE_AR_2SHOT_MAX_BYTES` (default `-1`): overrides the 2-shot AllReduce size cap from the arch table. Set to a positive value to override `ceNonRegMax[AR]`. Renamed from `RCCL_CE_AR_MAX_MSG_BYTES`. +* `RCCL_CE_AR_2SHOT_MIN_BYTES` (default `-1`): overrides the minimum message size for the CE 2-shot AllReduce path. When `-1`, resolved from `ceNonRegMin[AllReduce]` in the per-arch table. +* `RCCL_CE_AR_REG_MAX_MSG_BYTES` (default `-1`): overrides the registered CE AllReduce size cap from the arch table. +* `RCCL_CE_AR_STAGING_BYTES` (default `-1`): overrides the CE AllReduce staging buffer allocation size; when unset, `NCCL_CE_AR_STAGING_BYTES` is used. +* `RCCL_CE_COLL_MAX_BYTES` (default `-1`): overrides the CE registered-window and 2-shot size cap for all non-AllReduce collectives (AllGather, AlltoAll, etc.). When `-1`, resolved from `ceRegMax[func]` or `ceNonRegMax[func]` in the per-arch table. +* `RCCL_CE_COLL_MIN_BYTES` (default `-1`): overrides the minimum message size for CE non-AllReduce paths (both registered-window and scratch). When `-1`, resolved from `ceNonRegMin[func]` in the per-arch table. +* `RCCL_SYM_K_MAX_BYTES` (default `-1`): overrides the symmetric-kernel upper bound for all collectives. When `-1`, resolved from `symMaxR2[func]` (eager) or `symMaxR2Graph[func]` (graph-capture) in the per-arch table. `kThreshUnlimited` (SIZE_MAX) in the table means no suppression. +* `RCCL_SYM_K_MIN_BYTES` (default `-1`): overrides the symmetric-kernel lower bound for all collectives. When `-1`, resolved from `symMinR2[func]` in the per-arch table. Below this size, DDA is preferred over symk. +* Per-architecture dispatch table (`rcclArchThresholds`) centralizing algo/proto selection for DDA protocol, CE, and symmetric kernel per collective, replacing scattered hardcoded defaults. Tables are defined for gfx1250, gfx950, and gfx942; the lookup `rcclGetArchThresholds()` maps GCN arch strings to the appropriate entry. The `symMinR2[func]` field sets a per-collective lower bound for registered buffers. +* `RCCL_IGNORE_ARCH_TABLE` (default `0`): when set to `1`, bypasses the `rcclArchThresholds` dispatch table and falls back to compile-time constants for all DDA, CE, and symmetric-kernel thresholds. Has no effect on gfx942 and gfx950, which always use the arch table. +* New `rcclAddonAlgos_t` values: `RCCL_CE_SCRATCH` (CE via DDA scratch, distinct from `RCCL_CE_REGISTERED`), `RCCL_A2A_PIVOT`, `RCCL_A2A_GDA`, `RCCL_A2A_GIN_SDMA`, and `RCCL_DIRECT_ALLTOALL`, with corresponding `rcclGetAlgoName()` labels. +* rccl-tests: AlltoAll now reports algo/protocol in VERSION output, consistent with other collectives. +* Compatibility with NCCL 2.31.2. +* `net_ib_cast` refreshes local GIDs after `IBV_EVENT_GID_CHANGE` and applies them to data, flush, probing and port recovery QPs during IB port recovery, matching `net_ib` in NCCL 2.31.2. +* Per-collective configuration APIs (`ncclCollConfig_t` / `nccl*Config()` entry points) and the `ncclConfigExt_t` vendor extension list. Initialize configs with `NCCL_COLLCONFIG_INITIALIZER`. +* Communicator config (`ncclConfig_v23100`) fields for implicit launch ordering (`launchOrderImplicit`), RMA signal count (`numRmaSig`), eager RMA init (`rmaEagerInit`), and host collective fault tolerance (`hostCftMode`). +* NCCL profiler plugin API v7, with per-call user profiler tags and symmetric-kernel phase events. +* RAS diagnostics (`NCCL_RUN_RAS_DIAGNOSTICS`) covering GPU inventory, ROCm runtime versions, ECC counters, XGMI link state, and `NCCL_*` environment consistency (AMDSMI in place of NVML). +* Communicator init diagnostics (`NCCL_RUN_DIAGNOSTICS`) with an active P2P connectivity check. +* Multiple GIN proxy progress threads via `NCCL_GIN_PROXY_NTHREADS`. +* Enabled hierarchical Copy Engine `ncclAllGather`, `ncclAlltoAll`, and `ncclAlltoAllv` on multi-node communicators: the intra-node phase runs over SDMA/CE across the LSA team and the inter-node phase over one-sided RMA through the CPU (GIN) proxy. The path is off by default and takes two opt-ins: `NCCL_CTA_POLICY=2` (equivalently `config.CTAPolicy = NCCL_CTA_POLICY_ZERO`), and symmetric registration of both the send and the receive buffer through `ncclCommWindowRegister`. On architectures other than gfx1250 (MI450) a third is needed, `NCCL_CUMEM_ENABLE=1`, because symmetric memory depends on cuMem and cuMem auto-detection enables itself only on gfx1250. `NCCL_GIN_ENABLE`, `NCCL_NUM_RMA_CTX`, `NCCL_WIN_ENABLE` and `NCCL_DMABUF_ENABLE` need no change, since their defaults already qualify; the inter-node rail draws from the internal RMA context pool (`NCCL_NUM_RMA_INT_CTX`, default 4) rather than the user-addressable contexts, so `numRmaCtx` only has to stay above zero. Selection also requires an available GIN backend and the same number of ranks on every node; communicators that do not qualify keep using the kernel path. `ncclAlltoAllv` walks the gathered per-peer size matrix, skips zero-length peers, and chunks each inter-node put; single-node AlltoAllv stays on the existing CE path. Validated for correctness and path selection from 1 to 32 nodes. +* Added `ncclGinFenceLevel` barrier semantics on the Anvil SDMA GIN backend (`NCCL_GIN_TYPE=7`). `gin.get`, `flushAsync` and `wait` are implemented instead of trapping, and a barrier's signal is now ordered behind the payload on the peer's SDMA queue, so `Put`, `Get` and the default `Put | Get` fences drain what they promise. `Put` covers a rank's puts to itself, and `ncclGinAllContexts` fences every GIN context. Single-node validated on gfx950. +* Added gfx1250 (MI450) Tensor Data Mover staging to the symmetric AllGather, AllReduce and ReduceScatter kernels, where NCCL uses NVIDIA's TMA. Their deep loops move each tile through the warp's shared-memory scratch with the DMA engine instead of per-lane vector loads; on ROCm that path previously compiled out entirely, so the `Tma` kernels were never generated. Launch tuning was also corrected so these kernels can use up to a quarter of a GPU's WGPs, where only 14 were reached before, which is where most of the large-message gain comes from. Reached only when the application opts into symmetric memory; `NCCL_SYM_TMA_ENABLE` now defaults to `1` (see Changed in RCCL 2.32.3). `NCCL_SYM_TMA_ENABLE=2` additionally forces these kernels wherever the collective has one, skipping both the tuner's cost comparison and the deep-loop size bar, which is intended for A/B measurement rather than production. The new kernels are generated only when gfx1250 is in `GPU_TARGETS`, and every other architecture keeps the vector path unchanged. +* Removed GIN rocSHMEM GDA dependency on rocSHMEM GDA bitcode: GIN rocSHMEM GDA is now header-only. The CMake `roc::rccl` target provides the required include directories in its `INSTALL_INTERFACE` property. + +### Changed + * Vendored `{fmt}` 10.2.1 as a header-only tree under `external/fmt`. RCCL, its tests, and the `topo_expl` tool now always build against this copy instead of looking for a system `fmt` package or downloading one through `FetchContent` or `git clone`. * `rcclSelectAllGather` signature changed — it now takes a `cudaStream_t stream` parameter (for live dispatch) and a `bool query` mode that uses `graphCapturingHint` instead of probing the stream. The comment block was updated too. * `RCCL_DDA_THRESHOLD`, `RCCL_DDA_LL_THRESHOLD`, and `RCCL_DDA_LL128_THRESHOLD` now default to `-1` (unset), resolved at runtime from the arch dispatch table instead of being hardcoded to 128 MiB and 64 KiB respectively (RCCL_DDA_LL128_THRESHOLD had no pre-table default; the LL128 tier was off by default). Setting these env vars explicitly still takes precedence over the table. diff --git a/projects/rccl/cmake/DeviceLinker.cmake b/projects/rccl/cmake/DeviceLinker.cmake index 185f090cfdc7..91be4ba25e4d 100644 --- a/projects/rccl/cmake/DeviceLinker.cmake +++ b/projects/rccl/cmake/DeviceLinker.cmake @@ -925,6 +925,7 @@ endif() # 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) @@ -1309,13 +1310,13 @@ if(GENERATE_SYM_KERNELS) # .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(_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) diff --git a/projects/rccl/docs/api-reference/env-variables.rst b/projects/rccl/docs/api-reference/env-variables.rst index b973c95046d5..b1c9f0bf23e1 100644 --- a/projects/rccl/docs/api-reference/env-variables.rst +++ b/projects/rccl/docs/api-reference/env-variables.rst @@ -775,6 +775,19 @@ 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 + | 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. + - | ``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. @@ -1300,4 +1313,3 @@ threshold table introduced for gfx1250. | ``2``: Forced. Wherever the collective has a TDM kernel, it is selected regardless of predicted time or message size. Intended for A/B measurement, not production. - diff --git a/projects/rccl/src/CMakeLists.txt b/projects/rccl/src/CMakeLists.txt index e4373bd21166..9d6de58defbb 100644 --- a/projects/rccl/src/CMakeLists.txt +++ b/projects/rccl/src/CMakeLists.txt @@ -169,6 +169,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 53152ed74158..5adeefaf1b5a 100755 --- a/projects/rccl/src/algorithms/gin/gin_all_reduce.h +++ b/projects/rccl/src/algorithms/gin/gin_all_reduce.h @@ -2,10 +2,12 @@ * 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 GinScatter (64 MiB to 4 GiB) is opt-in via + * 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, 256 MiB)). + * LSA two-shot ((4 MiB, GIN floor)). * See LICENSE.txt for license information. ******************************************************************************/ @@ -19,12 +21,16 @@ #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 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. @@ -36,10 +42,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 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; uint32_t* intraGpuCtaBar; + void* scratch; + size_t scratchBytes; + struct ncclDevrWindow* scratchWin; }; #if defined(ENABLE_ROCSHMEM_GIN) @@ -50,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 < 256 MiB 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 8daf4d49541d..ad20150c95f6 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 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. ******************************************************************************/ @@ -22,9 +28,150 @@ 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. +// 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; + 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, 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, + 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; +} + +// 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 // per-rank slice so vector loads stay aligned. inline bool ginAllReduceTwoShotEligible(size_t count, size_t typeSize, int nRanks) { @@ -40,26 +187,48 @@ 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 || bytes > kGinAllReduceGinScatterMaxBytes) return false; + if (nRanks <= 0) return false; + if (scratchBytes >= ginAllReduceGinScatterStagedScratchBytes(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. +// 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; - return count * typeSize < kGinAllReduceGinTwoShotMinBytes; + return count * typeSize < ginMinBytes; +} + +// 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 && 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 514a38cc004e..22dc019d347f 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,12 @@ * 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 + * 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. ******************************************************************************/ @@ -14,6 +20,8 @@ #include "nccl_device.h" #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; @@ -61,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 @@ -331,4 +312,229 @@ __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) { + 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; + } + gin.put(world, dst, scratchWin, destSlotOff, sendWin, + sendOff + static_cast(dst) * perRankBytes + sendChunkOff, chunkBytes, + ncclGin_WeakSignalInc{signalIndex}); + } +} + +__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) { + 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); +} + +// 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); + } +} + +// Thread 0 waits until every other rank's own signal has advanced by `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; + 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). 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; + } + 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 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) +#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; + + // 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}; + 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); + + // 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) { + 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); + 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) { + 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); + } + 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) { + 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); + __syncthreads(); + } + } + } else { + for (int c = 0; c < nChunks; ++c) { + if (blockIdx.x == 0) { + if (c > 0 && staged) { + 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); + ginScatterWaitPeersThread0(gin, kGinScatterSignalBase, scatterBase, devComm.rank, nRanks, + static_cast(c + 1)); + } + 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); + __syncthreads(); + } + } + } + + if (blockIdx.x == 0) { + // 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{kGinScatterAgSignal}); + } + ginScatterWaitThread0(gin, kGinScatterAgSignal, agValue + static_cast(nRanks - 1)); + __syncthreads(); + gin.flush(cta); + } +} + } // 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..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 @@ -2,12 +2,16 @@ * 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 GinScatter is opt-in + * (NCCL_GIN_ALLREDUCE_GFX1250_ENABLE=1) for messages in [64 MiB, 4 GiB]. + * 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, 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, 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. @@ -30,9 +34,19 @@ #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); +// 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 // 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 +55,85 @@ 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; + ncclResult_t ret = ncclSuccess; + if (state->scratchWin != nullptr) { + NCCLCHECKIGNORE(ncclCommWindowDeregister(comm, state->scratchWin->vidmem), ret); + state->scratchWin = nullptr; + } + if (state->scratch != nullptr) { + NCCLCHECKIGNORE(ncclMemFree(state->scratch), ret); + state->scratch = nullptr; + } + state->scratchBytes = 0; + return ret; +} + +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); + + // 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) { + 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); + } + } + } + + CUDACHECKIGNORE(cudaThreadExchangeStreamCaptureMode(&captureMode)); + return ret; +} + // 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: // @@ -65,6 +158,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; @@ -200,6 +294,44 @@ 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) { + // 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 = + 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; + } + + 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; +} + template static ncclResult_t ncclAllReduceGinSdmaTyped(const void* sendbuff, void* recvbuff, size_t count, ncclComm* comm, cudaStream_t stream, struct ncclDevrWindow* sendWin, @@ -210,6 +342,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 +360,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")) 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; @@ -252,14 +391,28 @@ 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); + 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, 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,13 +437,15 @@ ncclResult_t ncclAllReduceGinSdma(const void* sendbuff, void* recvbuff, size_t c ncclResult_t ncclGinAllReduceFinalize(ncclComm* comm) { struct ncclGinAllReduceState* state = &comm->ginAllReduceState; + 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 c5cc5c8f6993..cfff6d02aafc 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 7339322c3529..f6865ffcdd99 100644 --- a/projects/rccl/src/rccl_wrap.cc +++ b/projects/rccl/src/rccl_wrap.cc @@ -1883,10 +1883,13 @@ ncclResult_t rcclSelectAllReduce(struct ncclComm* comm, const void* sendbuff, vo // it does not hand 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 - // !symEligible gate because GIN already returned above for those sizes. + // 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). 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 8658b547566d..4016ea52634f 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -717,6 +717,33 @@ 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, 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 + ${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") + # utility.h as plain C++ for its !NCCL_DEVICE_COMPILE branches; not with CXX=hipcc, which defines __HIPCC__ and fails strict CI. add_executable(rccl-UnitTestsUtilityHostOnly NcclUtilityHostOnlyTests.cpp) set_source_files_properties(NcclUtilityHostOnlyTests.cpp PROPERTIES LANGUAGE CXX) 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) 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..3a8c7bcb267f 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" @@ -27,8 +28,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, 64ull * kMiB); + EXPECT_EQ(kGinAllReduceGinScatterNumSlots, 2); } TEST(GinAllReducePolicy, DefaultRejectsBelow256MiB) { @@ -107,4 +114,215 @@ 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) { + // 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) { + 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)); +} + +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)); + EXPECT_FALSE(ginAllReduceGinScatterLaunch(kGinAllReduceGinScatterMaxBytes + 16, 512ull * kMiB, kRanks, scratch)); + EXPECT_FALSE(ginAllReduceGinScatterLaunch(64ull * kMiB, chunk64, kRanks, 0)); +} + +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 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, 8); + 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, 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, 1); + 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, 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(reinterpret_cast(staged), + reinterpret_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); + 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); +} + } // namespace 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..cc02e9d3e900 --- /dev/null +++ b/projects/rccl/test/gin/gin_scatter_reduce_gpu_test.cpp @@ -0,0 +1,277 @@ +/************************************************************************* + * 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, 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 + +#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, 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; + // 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); + 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 << " strideElems=" << strideElems; + + HIP_OK(hipFree(dSelf)); + HIP_OK(hipFree(dIncoming)); + 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) { + 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); + // 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); +} + +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 7151e1214506..765bb9c28126 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 @@ -2090,6 +2090,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", @@ -3998,6 +4013,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 5fb99b94ef68..79d2ac6c7ca9 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 @@ -1620,6 +1620,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", @@ -5573,6 +5588,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 c1c6b5748e65..091f88b81499 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 @@ -1715,6 +1715,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", @@ -5768,6 +5783,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 1097725b8e8d..5ffd033e4290 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 @@ -2158,6 +2158,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", @@ -4238,6 +4253,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 b66da890dc41..c6239691f049 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 @@ -2327,6 +2327,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", @@ -7180,6 +7195,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 bdfde85a6385..5e44c6e8198d 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 @@ -2843,6 +2843,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", @@ -6231,6 +6246,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",