feat(rccl): GIN SDMA alltoall

projects/rccl/cmake/DeviceLinker.cmake

@@ -351,7 +351,9 @@ foreach(DL_GPU_TARGET ${DL_GPU_TARGETS})
   # present in the device ELF — they cannot be imported from a shared library.
   set(_rocshmem_bitcode_arg "")
   set(_rocshmem_link_depends "")
-  if((ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN) AND ROCSHMEM_INSTALL_DIR)
+  # Not ENABLE_ROCSHMEM_GIN: no object here references rocSHMEM, and the full
+  # device bitcode carries unresolved IPC and RO symbols that break the ELF load.
+  if(ENABLE_ROCSHMEM AND ROCSHMEM_INSTALL_DIR)
     set(_rocshmem_bc "${ROCSHMEM_INSTALL_DIR}/lib/librocshmem_device_${DL_GPU_TARGET}.bc")
     set(_rocshmem_bitcode_arg "--rocshmem-bitcode=${_rocshmem_bc}")
     if(TARGET rocshmem_static)
@@ -665,6 +667,39 @@ add_custom_command(
   VERBATIM
 )
 
+# ===========================================================================
+# gin_alltoall_sdma.cu.cpp: GIN-SDMA alltoall 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_ALLTOALL_SDMA_FAT_OBJ "")
+if(ENABLE_ROCSHMEM_GIN)
+  set(GIN_ALLTOALL_SDMA_FAT_OBJ "${DEVICE_BUILD_DIR}/gin_alltoall_sdma.o")
+  add_custom_command(
+    OUTPUT  ${GIN_ALLTOALL_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_ALLTOALL_SDMA_FAT_OBJ}
+      ${HIPIFY_DIR}/src/algorithms/gin/sdma/gin_alltoall_sdma.cu.cpp
+    DEPENDS ${HIPIFY_DIR}/src/algorithms/gin/sdma/gin_alltoall_sdma.cu.cpp
+    COMMENT "DL compile: gin_alltoall_sdma.cu.cpp (GIN-SDMA alltoall kernel)"
+    VERBATIM
+  )
+endif()
+
 # ===========================================================================
 # dda_all_reduce_fabric.cu.cpp: fabric/VMM counterpart of the IPC file above.
 # ===========================================================================
@@ -1069,9 +1104,14 @@ 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}
+  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}
 )
 add_dependencies(device_linker_build hipify_all copy_nccl_device_headers)
+if((ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN) AND TARGET rocshmem_static)
+  # The fat objects above include GIN device headers, which pull in rocSHMEM
+  # headers installed to ext/rocshmem/include by the ExternalProject.
+  add_dependencies(device_linker_build rocshmem_static)
+endif()
 
 set(DEVICE_LINKER_OBJECTS
   ${COMMON_FAT_OBJ}
@@ -1095,6 +1135,7 @@ set(DEVICE_LINKER_OBJECTS
   ${DDA_REDUCE_SCATTER_FABRIC_LL_FAT_OBJ}
   ${DDA_REDUCE_SCATTER_FABRIC_LL128_FAT_OBJ}
   ${SYM_FAT_OBJS}
+  ${GIN_ALLTOALL_SDMA_FAT_OBJ}
 )
 
 # ===========================================================================

projects/rccl/src/CMakeLists.txt

@@ -139,6 +139,7 @@ set(SRC_FILES
   algorithms/dda/alltoall/alltoall_dda_fabric_ll128.h
   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
   include/alltoallv_meta.h
   include/alt_rsmi.h
   include/archinfo.h
@@ -520,6 +521,7 @@ if(ENABLE_ROCSHMEM_GIN)
     src/gin/gin_rocshmem_gda_factory.cc
     src/gin/gin_plugin_anvil_sdma.cc
     src/gin/gin_anvil_ipc_table_host.cc
+    src/algorithms/gin/sdma/gin_alltoall_sdma.cu
   )
 endif()
 
@@ -812,6 +814,9 @@ if(ENABLE_DEVICE_LINKER)
   list(FILTER HIP_SOURCES EXCLUDE REGEX "dda_reduce_scatter_ipc\\.cu\\.cpp$")
   list(FILTER HIP_SOURCES EXCLUDE REGEX "dda_all_gather_ipc\\.cu\\.cpp$")
   list(FILTER HIP_SOURCES EXCLUDE REGEX "dda_alltoall_ipc\\.cu\\.cpp$")
+  # 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 "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$")
@@ -1266,12 +1271,22 @@ endif()
 # a FetchContent-built fmt-header-only target leaks into rccl's link
 # interface and breaks install(EXPORT) for static builds.
 target_link_libraries(rccl PRIVATE $<BUILD_INTERFACE:fmt::fmt-header-only>)
+if(ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN)
+  # Resolve to a full path: nothing on the link line supplies ${ROCM_PATH}/lib, so a bare
+  # -lhsakmt would have no search path to resolve against.
+  find_library(HSAKMT_LIBRARY NAMES hsakmt
+    HINTS "${ROCM_PATH}"
+    PATH_SUFFIXES lib lib64)
+  if(NOT HSAKMT_LIBRARY)
+    message(FATAL_ERROR "Could not find libhsakmt")
+  endif()
+endif()
 if(ENABLE_ROCSHMEM)
   # librocshmem.a is compiled with -fgpu-rdc (relocatable device code).
   # --hip-link tells amdclang++ to perform the RDC device-link step as part
   # of the final librccl.so link.
   target_link_libraries(rccl PRIVATE ${ROCSHMEM_LIBRARY})
-  target_link_libraries(rccl PRIVATE hsakmt)
+  target_link_libraries(rccl PRIVATE ${HSAKMT_LIBRARY})
   target_link_libraries(rccl PRIVATE ${IBVERBS})
   set(_rocshmem_offload_arch_flags "")
   foreach(_gpu_raw ${GPU_TARGETS})
@@ -1284,7 +1299,55 @@ elseif(ENABLE_ROCSHMEM_GIN)
   # GIN device templates (Put/Flush/etc.) reference rocshmem device symbols.
   # With the device linker (default), per-arch .bc files are injected directly
   # into the per-arch device link step via DeviceLinker.cmake.
-  # Host symbols from GIN plugins are left unresolved via --allow-shlib-undefined.
+  #
+  # Pull the SDMA host objects out of librocshmem.a so apps running only SDMA do
+  # not have to link it. Seeding by symbol rather than by file survives renames.
+  # GDA stays unresolved, as it needs constmem registration this mode cannot do.
+  set(_gin_rocshmem_syms
+    gin_anvil_sdma_probe
+    gin_anvil_sdma_create
+    gin_anvil_sdma_destroy
+    gin_anvil_sdma_get_num_channels
+    gin_anvil_sdma_get_channel_stride
+    _ZN8rocshmem3ibvE)
+  set(_gin_rocshmem_undefine "")
+  foreach(_sym ${_gin_rocshmem_syms})
+    list(APPEND _gin_rocshmem_undefine -u ${_sym})
+  endforeach()
+
+  set(_gin_rocshmem_object "${CMAKE_CURRENT_BINARY_DIR}/rocshmem_host.o")
+  set(_gin_rocshmem_archive_name "librocshmem_gin_host.a")
+  set(_gin_rocshmem_archive "${CMAKE_CURRENT_BINARY_DIR}/${_gin_rocshmem_archive_name}")
+  if(TARGET rocshmem_static)
+    set(_gin_rocshmem_depends rocshmem_static)
+  else()
+    set(_gin_rocshmem_depends ${ROCSHMEM_LIBRARY})
+  endif()
+
+  add_custom_command(
+    OUTPUT  ${_gin_rocshmem_object}
+    COMMAND ${CMAKE_LINKER} -r ${_gin_rocshmem_undefine} ${ROCSHMEM_LIBRARY} -o ${_gin_rocshmem_object}
+    DEPENDS ${_gin_rocshmem_depends}
+    COMMENT "Extracting GIN host objects from librocshmem.a"
+    VERBATIM)
+
+  # An app linking rocshmem under -rdynamic would interpose our globals, so both
+  # modules would run destructors over one object and double free at exit.
+  # --exclude-libs keeps our copies local. It needs an archive, hence the wrapper.
+  add_custom_command(
+    OUTPUT  ${_gin_rocshmem_archive}
+    COMMAND ${CMAKE_COMMAND} -E rm -f ${_gin_rocshmem_archive}
+    COMMAND ${CMAKE_AR} rcs ${_gin_rocshmem_archive} ${_gin_rocshmem_object}
+    DEPENDS ${_gin_rocshmem_object}
+    COMMENT "Packaging GIN host objects into ${_gin_rocshmem_archive_name}"
+    VERBATIM)
+  add_custom_target(gin_rocshmem_host_objects DEPENDS ${_gin_rocshmem_archive})
+
+  add_dependencies(rccl gin_rocshmem_host_objects)
+  target_link_libraries(rccl PRIVATE ${_gin_rocshmem_archive})
+  target_link_options(rccl PRIVATE "LINKER:--exclude-libs,${_gin_rocshmem_archive_name}")
+  # anvil.cpp calls hsaKmt* directly. hsa-runtime64 is already linked above.
+  target_link_libraries(rccl PRIVATE ${HSAKMT_LIBRARY})
   target_link_options(rccl PRIVATE "LINKER:--allow-shlib-undefined")
 endif()
 

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

@@ -0,0 +1,44 @@
+/*************************************************************************
+ * Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * Host entry points for the GIN-SDMA all-to-all path launched from ncclAlltoAll.
+ * See LICENSE.txt for license information.
+ ************************************************************************/
+
+#ifndef GIN_ALLTOALL_H_
+#define GIN_ALLTOALL_H_
+
+#include "nccl.h"
+#include "nccl_device.h"
+
+struct ncclComm;
+
+// Upper bound on CTAs, and so on the LSA barriers and GIN signals reserved here.
+constexpr int kGinA2AMaxCtas = 64;
+
+// Lazily created on the first eligible alltoall and torn down with the comm.
+// Declared unconditionally, or rcclras and librccl disagree on ncclComm layout.
+struct ncclGinA2AState {
+  bool initialized;
+  struct ncclDevComm devComm;
+};
+
+#if defined(ENABLE_ROCSHMEM_GIN)
+
+bool ncclAllToAllGinSdmaEligible(ncclComm* comm, const void* sendbuff, void* recvbuff, size_t count,
+                                 ncclDataType_t datatype);
+
+ncclResult_t ncclAllToAllGinSdma(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype,
+                                 ncclComm* comm, cudaStream_t stream);
+
+ncclResult_t ncclGinA2AFinalize(ncclComm* comm);
+
+#else
+
+// gin_alltoall_sdma.cu is only compiled with the GIN backend, so teardown has
+// nothing to release here.
+inline ncclResult_t ncclGinA2AFinalize(ncclComm*) { return ncclSuccess; }
+
+#endif
+
+#endif