feat(rccl): Scaleup Allreduce using GIN/LSA

projects/rccl/cmake/DeviceLinker.cmake

@@ -700,6 +700,38 @@ if(ENABLE_ROCSHMEM_GIN)
   )
 endif()
 
+# ===========================================================================
+# gin_all_reduce_sdma.cu.cpp: GIN-SDMA allreduce kernel.
+#
+# The defines go here because this file is filtered out of the rccl target, so
+# per-source properties never apply. Leaving only SDMA on keeps ncclGinCallImpl
+# on its single backend branch. GDA must stay off, as this object gets no QP bitcode.
+# ===========================================================================
+set(GIN_ALLREDUCE_SDMA_FAT_OBJ "")
+if(ENABLE_ROCSHMEM_GIN)
+  set(GIN_ALLREDUCE_SDMA_FAT_OBJ "${DEVICE_BUILD_DIR}/gin_all_reduce_sdma.o")
+  add_custom_command(
+    OUTPUT  ${GIN_ALLREDUCE_SDMA_FAT_OBJ}
+    COMMAND ${DL_CLANG}
+      -x hip ${DL_OFFLOAD_ARCH_FLAGS}
+      ${DL_HIP_COMPILER_FLAGS}
+      -DRCCL_DEVICE_LINKER
+      -DNCCL_GIN_ANVIL_SDMA_ENABLE=1
+      -DNCCL_GIN_PROXY_ENABLE=0
+      -DNCCL_GIN_ROCSHMEM_GDA_ENABLE=0
+      ${_link_def_flags}
+      ${_host_inc_flags}
+      ${DL_OPT_FLAGS}
+      ${DL_INHERITED_FLAGS}
+      -std=c++17
+      -fPIC
+      -c -o ${GIN_ALLREDUCE_SDMA_FAT_OBJ}
+      ${HIPIFY_DIR}/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu.cpp
+    DEPENDS ${HIPIFY_DIR}/src/algorithms/gin/sdma/gin_all_reduce_sdma.cu.cpp
+    COMMENT "DL compile: gin_all_reduce_sdma.cu.cpp (GIN-SDMA allreduce kernel)"
+    VERBATIM
+  )
+endif()
 # ===========================================================================
 # dda_all_reduce_fabric.cu.cpp: fabric/VMM counterpart of the IPC file above.
 # ===========================================================================
@@ -1104,7 +1136,7 @@ endif()
 # Top-level target
 # ===========================================================================
 add_custom_target(device_linker_build ALL
-  DEPENDS ${COMMON_FAT_OBJ} ${ONERANK_FAT_OBJ} ${COLLECTIVES_FAT_OBJ} ${DDA_ALL_REDUCE_IPC_FAT_OBJ} ${DDA_REDUCE_SCATTER_IPC_FAT_OBJ} ${DDA_ALL_GATHER_IPC_FAT_OBJ} ${DDA_ALLTOALL_IPC_FAT_OBJ} ${DDA_ALL_REDUCE_FABRIC_FAT_OBJ} ${DDA_ALL_REDUCE_FABRIC_LL_FAT_OBJ} ${DDA_ALL_REDUCE_FABRIC_LL128_FAT_OBJ} ${DDA_REDUCE_SCATTER_FABRIC_FAT_OBJ} ${DDA_ALL_GATHER_FABRIC_FAT_OBJ} ${DDA_ALL_GATHER_FABRIC_LL_FAT_OBJ} ${DDA_ALL_GATHER_FABRIC_LL128_FAT_OBJ} ${DDA_ALLTOALL_FABRIC_FAT_OBJ} ${DDA_ALLTOALL_FABRIC_LL_FAT_OBJ} ${DDA_ALLTOALL_FABRIC_LL128_FAT_OBJ} ${DDA_REDUCE_SCATTER_FABRIC_LL_FAT_OBJ} ${DDA_REDUCE_SCATTER_FABRIC_LL128_FAT_OBJ} ${CE_REDUCE_FAT_OBJS} ${SYM_FAT_OBJS} ${GIN_ALLTOALL_SDMA_FAT_OBJ}
+	DEPENDS ${COMMON_FAT_OBJ} ${ONERANK_FAT_OBJ} ${COLLECTIVES_FAT_OBJ} ${DDA_ALL_REDUCE_IPC_FAT_OBJ} ${DDA_REDUCE_SCATTER_IPC_FAT_OBJ} ${DDA_ALL_GATHER_IPC_FAT_OBJ} ${DDA_ALLTOALL_IPC_FAT_OBJ} ${DDA_ALL_REDUCE_FABRIC_FAT_OBJ} ${DDA_ALL_REDUCE_FABRIC_LL_FAT_OBJ} ${DDA_ALL_REDUCE_FABRIC_LL128_FAT_OBJ} ${DDA_REDUCE_SCATTER_FABRIC_FAT_OBJ} ${DDA_ALL_GATHER_FABRIC_FAT_OBJ} ${DDA_ALL_GATHER_FABRIC_LL_FAT_OBJ} ${DDA_ALL_GATHER_FABRIC_LL128_FAT_OBJ} ${DDA_ALLTOALL_FABRIC_FAT_OBJ} ${DDA_ALLTOALL_FABRIC_LL_FAT_OBJ} ${DDA_ALLTOALL_FABRIC_LL128_FAT_OBJ} ${DDA_REDUCE_SCATTER_FABRIC_LL_FAT_OBJ} ${DDA_REDUCE_SCATTER_FABRIC_LL128_FAT_OBJ} ${CE_REDUCE_FAT_OBJS} ${SYM_FAT_OBJS} ${GIN_ALLTOALL_SDMA_FAT_OBJ} ${GIN_ALLREDUCE_SDMA_FAT_OBJ}
 )
 add_dependencies(device_linker_build hipify_all copy_nccl_device_headers)
 if((ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN) AND TARGET rocshmem_static)
@@ -1136,6 +1168,7 @@ set(DEVICE_LINKER_OBJECTS
   ${DDA_REDUCE_SCATTER_FABRIC_LL128_FAT_OBJ}
   ${SYM_FAT_OBJS}
   ${GIN_ALLTOALL_SDMA_FAT_OBJ}
+  ${GIN_ALLREDUCE_SDMA_FAT_OBJ}
 )
 
 # ===========================================================================

projects/rccl/src/CMakeLists.txt

@@ -140,6 +140,9 @@ set(SRC_FILES
   algorithms/dda/reduce_scatter/reduce_scatter_dda_fabric_ll.h
   algorithms/dda/reduce_scatter/reduce_scatter_dda_fabric_ll128.h
   algorithms/gin/gin_alltoall.h
+  algorithms/gin/gin_all_reduce.h
+  algorithms/gin/gin_all_reduce_policy.h
+  algorithms/gin/sdma/all_reduce_gin_sdma.h 
   include/alltoallv_meta.h
   include/alt_rsmi.h
   include/archinfo.h
@@ -522,6 +525,7 @@ if(ENABLE_ROCSHMEM_GIN)
     src/gin/gin_plugin_anvil_sdma.cc
     src/gin/gin_anvil_ipc_table_host.cc
     src/algorithms/gin/sdma/gin_alltoall_sdma.cu
+    src/algorithms/gin/sdma/gin_all_reduce_sdma.cu 
   )
 endif()
 
@@ -817,6 +821,7 @@ if(ENABLE_DEVICE_LINKER)
   # Contains the GIN-SDMA alltoall __global__ kernel and its launch site, so it
   # needs full HIP compilation rather than --offload-host-only.
   list(FILTER HIP_SOURCES EXCLUDE REGEX "gin_alltoall_sdma\\.cu\\.cpp$")
+  list(FILTER HIP_SOURCES EXCLUDE REGEX "gin_all_reduce_sdma\\.cu\\.cpp$")
   list(FILTER HIP_SOURCES EXCLUDE REGEX "dda_all_reduce_fabric\\.cu\\.cpp$")
   list(FILTER HIP_SOURCES EXCLUDE REGEX "dda_all_reduce_fabric_ll\\.cu\\.cpp$")
   list(FILTER HIP_SOURCES EXCLUDE REGEX "dda_all_reduce_fabric_ll128\\.cu\\.cpp$")

projects/rccl/src/algorithms/gin/gin_all_reduce.h

@@ -0,0 +1,69 @@
+/*************************************************************************
+ * 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.
+ * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables LSA one-shot (<= 4 MiB) and
+ * LSA two-shot ((4 MiB, 256 MiB)).
+ * See LICENSE.txt for license information.
+ ******************************************************************************/
+
+#ifndef GIN_ALL_REDUCE_H_
+#define GIN_ALL_REDUCE_H_
+
+#include "gin_all_reduce_policy.h"
+#include "nccl.h"
+#include "nccl_device.h"
+
+#include <stdint.h>
+
+struct ncclComm;
+
+// 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).
+
+// 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.
+//
+// Deliberately holds no per-launch host epoch. All cross-rank synchronization lives on the
+// device (LSA barrier epochs in the resource window, GIN signals and their shadows) and is
+// re-read by the kernel on every launch, which is what lets these collectives be captured
+// into a graph and replayed: a host-side counter baked into a kernel argument at capture time
+// 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.
+struct ncclGinAllReduceState {
+  bool initialized;
+  struct ncclDevComm devComm;
+  uint32_t* intraGpuCtaBar;
+};
+
+#if defined(ENABLE_ROCSHMEM_GIN)
+
+// RCCL_GIN_ALLREDUCE_FORCE_ENABLE. Defined in gin_all_reduce_sdma.cu.
+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).
+bool ncclAllReduceGinSdmaYieldToDda(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count,
+                                    ncclDataType_t datatype, ncclRedOp_t op);
+
+ncclResult_t ncclAllReduceGinSdma(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype,
+                                  ncclRedOp_t op, ncclComm* comm, cudaStream_t stream);
+
+ncclResult_t ncclGinAllReduceFinalize(ncclComm* comm);
+
+#else
+
+inline ncclResult_t ncclGinAllReduceFinalize(ncclComm*) { return ncclSuccess; }
+
+#endif
+
+#endif

projects/rccl/src/algorithms/gin/gin_all_reduce_policy.h

@@ -0,0 +1,65 @@
+/*************************************************************************
+ * 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.
+ * See LICENSE.txt for license information.
+ ******************************************************************************/
+
+#ifndef GIN_ALL_REDUCE_POLICY_H_
+#define GIN_ALL_REDUCE_POLICY_H_
+
+#include <cstddef>
+
+constexpr int kGinAllReduceLsaCtas = 56;
+constexpr int kGinAllReduceLsaTwoShotCtasPerPeer = 8;
+constexpr int kGinAllReduceLsaTwoShotMaxCtas = kGinAllReduceLsaTwoShotCtasPerPeer * 16;
+constexpr int kGinAllReduceMaxRanks = 8;
+
+constexpr int kGinAllReduceMinBytes = 512ULL * 1024;
+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;
+
+constexpr size_t kGinAllReduceMinPutBytes = 128;
+
+// 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) {
+  if (nRanks <= 0 || typeSize == 0) return false;
+  if (count % static_cast<size_t>(nRanks) != 0) return false;
+  const size_t countPerRank = count / static_cast<size_t>(nRanks);
+  return (countPerRank * typeSize) % 16 == 0;
+}
+
+inline bool ginAllReduceGinTwoShotEligible(size_t count, size_t typeSize, int nRanks) {
+  if (!ginAllReduceTwoShotEligible(count, typeSize, nRanks)) return false;
+  const size_t chunkBytes = (count / static_cast<size_t>(nRanks)) * typeSize;
+  return chunkBytes >= kGinAllReduceMinPutBytes;
+}
+
+// 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) {
+  const size_t bytes = count * typeSize;
+  if (bytes < kGinAllReduceGinTwoShotMinBytes && !forceEnable) return false;
+  if (bytes < static_cast<size_t>(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) {
+    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) {
+  if (forceEnable) return false;
+  return count * typeSize < kGinAllReduceGinTwoShotMinBytes;
+}
+
+#endif