nusislam · Nov 2, 2026 · Oct 8, 2026 · Oct 9, 2026 · Oct 13, 2026 · Oct 22, 2026
@@ -22,6 +22,8 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * `NCCL_HIER_CE_COLL_AG_RAIL_RING_ENABLE` (default `-1`): a positive value selects a ring for the inter-node rail phase of hierarchical Copy Engine `ncclAllGather`. The default keeps the direct path.
 * `NCCL_IB_SORT_MERGE_NICS`: sorts the sub-devices of a merged IB device by plane ID. RCCL defaults it to `0` (NCCL defaults to `1`), so merged-device order and names are unchanged.
 * nccl4py: per-call collective configuration (`NCCLCollConfig`, `VendorOption`) including the launch completion event, communicator properties (`NCCLCommProperties`), and the `GIN_ONLY` window flag. On ROCm, the HIP `Event` shim provides only an event handle; `record()`, `sync()` and `query` are not implemented.
+* `NCCL_GIN_ALLREDUCE_GFX1250_ENABLE` (default `0`): opt in to GinScatter AllReduce on gfx1250 for symmetric sum messages from 64 MiB through 4 GiB. Left off, gfx1250 keeps the pre-existing AllReduce selectors.
+* `NCCL_GIN_ALLREDUCE_SCRATCH_BYTES` (default 512 MiB): GinScatter incoming scratch. `0` skips the allocation and GinScatter is not selected. The buffer is allocated only on the GinScatter path.
 * `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.
 
 ### Changed

@@ -915,12 +915,63 @@ if(ENABLE_ROCSHMEM_GIN)
   )
 endif()
 
+# ===========================================================================
+# rocshmem QP device bitcode for GIN GDA puts (QueuePair::put_nbi, etc.).
+# Used by GIN symmetric kernels. The installed per-arch .bc files are
+# arch-optimized and cannot be used with -mlink-builtin-bitcode across archs,
+# so this is unoptimized IR + constmem stubs.
+# ===========================================================================
+set(GIN_ROCSHMEM_QP_BC_FLAG "")
+set(GIN_ROCSHMEM_QP_DEPS "")
+if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR AND DL_GPU_TARGETS AND GENERATE_SYM_KERNELS)
+  list(GET DL_GPU_TARGETS 0 _bc_arch)
+  set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc")
+  find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED)
+  find_program(_llvm_dis  llvm-dis  HINTS ${ROCM_PATH}/llvm/bin REQUIRED)
+  find_program(_llvm_as   llvm-as   HINTS ${ROCM_PATH}/llvm/bin REQUIRED)
+
+  set(_cm_src "${CMAKE_SOURCE_DIR}/src/gin/gin_rocshmem_constmem.hip")
+  set(_cm_bc  "${DEVICE_BUILD_DIR}/gin_rocshmem_constmem.bc")
+  set(_qp_raw "${DEVICE_BUILD_DIR}/rocshmem_qp_raw.bc")
+
+  add_custom_command(
+    OUTPUT ${_cm_bc}
+    COMMAND ${DL_CLANG}
+      -x hip --cuda-device-only --offload-arch=${_bc_arch}
+      -emit-llvm -Xclang -disable-llvm-passes
+      -std=c++17 -fPIC
+      -I${ROCSHMEM_SOURCE_DIR}/src
+      -I${ROCSHMEM_SOURCE_DIR}/include
+      -c -o ${_cm_bc} ${_cm_src}
+    DEPENDS ${_cm_src}
+    COMMENT "DL: compiling rocshmem constmem stubs to device bitcode"
+    VERBATIM)
+
+  add_custom_command(
+    OUTPUT ${_qp_bc}
+    COMMAND ${_llvm_link} ${_cm_bc} -o ${_qp_raw}
+    COMMAND ${_llvm_dis} -o ${_qp_raw}.ll ${_qp_raw}
+    COMMAND grep -v -E "@llvm[.]compiler[.]used|@__hip_cuid_"
+      ${_qp_raw}.ll > ${_qp_raw}.clean.ll
+    COMMAND ${_llvm_as} ${_qp_raw}.clean.ll -o ${_qp_bc}
+    DEPENDS ${_cm_bc}
+    COMMENT "DL: linking rocshmem QP device bitcode (with constmem, stripped)"
+    VERBATIM)
+  add_custom_target(dl_rocshmem_qp_bc DEPENDS ${_qp_bc})
+  if(TARGET rocshmem_static)
+    add_dependencies(dl_rocshmem_qp_bc rocshmem_static)
+  endif()
+  set(GIN_ROCSHMEM_QP_BC_FLAG -Xclang -mlink-builtin-bitcode -Xclang ${_qp_bc})
+  set(GIN_ROCSHMEM_QP_DEPS dl_rocshmem_qp_bc)
+endif()
+
 # ===========================================================================
 # gin_all_reduce_sdma.cu.cpp: GIN-SDMA allreduce kernel.
 #
 # The defines go here because this file is filtered out of the rccl target, so
 # per-source properties never apply. Leaving only SDMA on keeps ncclGinCallImpl
 # on its single backend branch. GDA must stay off, as this object gets no QP bitcode.
+# gfx1250 GinScatter and gfx950 two-shot both use Anvil SDMA gin.put().
 # ===========================================================================
 set(GIN_ALLREDUCE_SDMA_FAT_OBJ "")
 if(ENABLE_ROCSHMEM_GIN)
@@ -1300,65 +1351,9 @@ endforeach()
 # ===========================================================================
 set(SYM_FAT_OBJS "")
 if(GENERATE_SYM_KERNELS)
-  # When ENABLE_ROCSHMEM_GIN is set, GIN device templates reference rocshmem
-  # device symbols (QueuePair::put_nbi, atomic_add, etc.). The installed per-arch
-  # .bc files are arch-optimized (opt -mcpu=) and can't be used with
-  # -mlink-builtin-bitcode across archs. Instead, llvm-link the pre-opt
-  # individual source .bc files (arch-agnostic) into a minimal QP-only bitcode.
-  set(_sym_rocshmem_bc_flag "")
-  set(_sym_rocshmem_deps "")
-  if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR)
-    # Pick the first arch's pre-opt bitcode dir (all archs produce identical
-    # unoptimized IR since -Xclang -disable-llvm-passes is used).
-    list(GET DL_GPU_TARGETS 0 _bc_arch)
-    set(_bc_dir "${ROCSHMEM_SOURCE_DIR}/build/bitcode/${_bc_arch}")
-    set(_qp_bc "${DEVICE_BUILD_DIR}/rocshmem_qp_device.bc")
-    find_program(_llvm_link llvm-link HINTS ${ROCM_PATH}/llvm/bin REQUIRED)
-    find_program(_llvm_dis  llvm-dis  HINTS ${ROCM_PATH}/llvm/bin REQUIRED)
-    find_program(_llvm_as   llvm-as   HINTS ${ROCM_PATH}/llvm/bin REQUIRED)
-
-    # Pipeline:
-    #  1. Compile gin_rocshmem_constmem.hip → device-only .bc (provides
-    #     rocshmem::constmem and rocshmem::logd_constants definitions that
-    #     queue_pair.bc references as external)
-    #  2. llvm-link QP .bc files + constmem .bc into one module
-    #  3. Strip @llvm.compiler.used and @__hip_cuid_ via text round-trip
-    #     (these AMDGCN addrspace(1) appending globals clash with the
-    #     host-side addrspace(0) equivalents in fat-object compilation)
-    set(_cm_src "${CMAKE_SOURCE_DIR}/src/gin/gin_rocshmem_constmem.hip")
-    set(_cm_bc  "${DEVICE_BUILD_DIR}/gin_rocshmem_constmem.bc")
-    set(_qp_raw "${DEVICE_BUILD_DIR}/rocshmem_qp_raw.bc")
-
-    add_custom_command(
-      OUTPUT ${_cm_bc}
-      COMMAND ${DL_CLANG}
-        -x hip --cuda-device-only --offload-arch=${_bc_arch}
-        -emit-llvm -Xclang -disable-llvm-passes
-        -std=c++17 -fPIC
-        -I${ROCSHMEM_SOURCE_DIR}/src
-        -I${ROCSHMEM_SOURCE_DIR}/include
-        -c -o ${_cm_bc} ${_cm_src}
-      DEPENDS ${_cm_src}
-      COMMENT "DL: compiling rocshmem constmem stubs to device bitcode"
-      VERBATIM)
-
-    add_custom_command(
-      OUTPUT ${_qp_bc}
-      COMMAND ${_llvm_link} ${_cm_bc} -o ${_qp_raw}
-      COMMAND ${_llvm_dis} -o ${_qp_raw}.ll ${_qp_raw}
-      COMMAND grep -v -E "@llvm[.]compiler[.]used|@__hip_cuid_"
-        ${_qp_raw}.ll > ${_qp_raw}.clean.ll
-      COMMAND ${_llvm_as} ${_qp_raw}.clean.ll -o ${_qp_bc}
-      DEPENDS ${_cm_bc}
-      COMMENT "DL: linking rocshmem QP device bitcode (with constmem, stripped)"
-      VERBATIM)
-    add_custom_target(dl_rocshmem_qp_bc DEPENDS ${_qp_bc})
-    if(TARGET rocshmem_static)
-      add_dependencies(dl_rocshmem_qp_bc rocshmem_static)
-    endif()
-    set(_sym_rocshmem_bc_flag -Xclang -mlink-builtin-bitcode -Xclang ${_qp_bc})
-    set(_sym_rocshmem_deps dl_rocshmem_qp_bc)
-  endif()
+  # GIN symmetric kernels reuse the QP bitcode generated above.
+  set(_sym_rocshmem_bc_flag "${GIN_ROCSHMEM_QP_BC_FLAG}")
+  set(_sym_rocshmem_deps "${GIN_ROCSHMEM_QP_DEPS}")
   file(GLOB _sym_srcs CONFIGURE_DEPENDS "${HIPIFY_DIR}/gensrc/symmetric/*.cpp")
   foreach(_sym_src IN LISTS _sym_srcs)
     get_filename_component(_sym_name "${_sym_src}" NAME_WE)

@@ -526,6 +526,17 @@ in the following table.
         | ``6``: rocSHMEM GDA (required for that backend to initialize)
         | ``7``: Anvil SDMA (also accepts unset)
 
+    * - | ``NCCL_GIN_ALLREDUCE_GFX1250_ENABLE``
+        | Opt in to GinScatter AllReduce on gfx1250.
+      - | ``0``: gfx1250 keeps the pre-existing AllReduce selectors (default)
+        | ``1``: symmetric ``ncclSum`` fp32/fp16/bf16 messages from 64 MiB
+          through 4 GiB use GinScatter when scratch can hold them
+
+    * - | ``NCCL_GIN_ALLREDUCE_SCRATCH_BYTES``
+        | GinScatter incoming scratch size, in bytes.
+      - | ``536870912`` (512 MiB) by default, matching the CE AllReduce pool
+        | ``0``: do not allocate scratch, and do not select GinScatter
+
     * - | ``NCCL_RMA_PLUGIN``
         | Selects external one-sided RMA plugins, which are also the backend
           the built-in GIN proxy forwards to.

@@ -164,6 +164,7 @@ set(SRC_FILES
   algorithms/gin/gin_all_reduce.h
   algorithms/gin/gin_all_reduce_policy.h
   algorithms/gin/sdma/all_reduce_gin_sdma.h
+  algorithms/gin/sdma/gin_scatter_reduce.h
   include/alloc.h
   include/allocator.h
   include/alltoallv_meta.h

@@ -2,10 +2,11 @@
  * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
  *
  * Host entry points for the GIN-SDMA AllReduce path selected by rcclSelectAllReduce
- * when symmetric windows are used. By default only messages >= 256 MiB
- * (GIN two-shot) take this path; smaller messages use DDA AllReduce.
+ * when symmetric windows are used. By default gfx950 takes GIN two-shot for
+ * messages >= 256 MiB. gfx1250 GinScatter (64 MiB to 4 GiB) is opt-in via
+ * NCCL_GIN_ALLREDUCE_GFX1250_ENABLE. Smaller messages use DDA AllReduce.
  * RCCL_GIN_ALLREDUCE_FORCE_ENABLE=1 also enables LSA one-shot (<= 4 MiB) and
- * LSA two-shot ((4 MiB, 256 MiB)).
+ * LSA two-shot ((4 MiB, GIN floor)).
  * See LICENSE.txt for license information.
  ******************************************************************************/
 
@@ -19,12 +20,16 @@
 #include <stdint.h>
 
 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 +41,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)
@@ -51,7 +62,7 @@ bool ncclAllReduceGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* re
                                   ncclDataType_t datatype, ncclRedOp_t op);
 
 // True when this AllReduce is a GIN-SDMA candidate that the default size policy
-// left for DDA (message < 256 MiB and RCCL_GIN_ALLREDUCE_FORCE_ENABLE is not 1).
+// left for DDA (message below the GIN floor and RCCL_GIN_ALLREDUCE_FORCE_ENABLE is not 1).
 bool ncclAllReduceGinSdmaYieldToDda(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count,
                                     ncclDataType_t datatype, ncclRedOp_t op);