refactor(rocshmem/gda): refactor QueuePair class

projects/rccl-tests/src/gin_rocshmem_gda_factory_test.cu

@@ -22,7 +22,7 @@
 #include <cassert>
 
 #include <gin/gin_rocshmem_gda_factory.h>
-#include <nccl_device/gin/rocshmem_gda/queue_pair_device.h>
+#include <nccl_device/gin/rocshmem_gda/gda/queue_pair_provider.hpp>
 
 #define HIP_CHECK(cmd) do {                                      \
   hipError_t e = (cmd);                                          \
@@ -54,7 +54,9 @@ __global__ void gin_put_kernel(rocshmem::QueuePair **qps,
     rocshmem::ActiveWFInfo wf_info(peer, rocshmem::ThreadScope::thread);
     printf("gin_put_kernel: peer=%d dst=%p src=%p bytes=%zu rkey=0x%x lkey=0x%x\n",
            peer, dst, src, nbytes, dst_rkey, src_lkey);
-    qps[peer]->put_nbi(dst, dst_rkey, src, src_lkey, nbytes, wf_info);
+    uintptr_t dst_raddr = reinterpret_cast<uintptr_t>(dst);
+    uintptr_t src_laddr = reinterpret_cast<uintptr_t>(src);
+    qps[peer]->put_nbi(dst_raddr, dst_rkey, src_laddr, src_lkey, nbytes, wf_info);
     printf("gin_put_kernel: put posted, calling quiet\n");
     qps[peer]->quiet(wf_info);
     printf("gin_put_kernel: quiet done\n");
@@ -69,7 +71,8 @@ __global__ void gin_atomic_kernel(rocshmem::QueuePair **qps,
     rocshmem::ActiveWFInfo wf_info(peer, rocshmem::ThreadScope::thread);
     printf("gin_atomic_kernel: peer=%d addr=%p rkey=0x%x val=%lld\n",
            peer, remote_addr, rkey, (long long)value);
-    qps[peer]->atomic_add(remote_addr, rkey, value, wf_info);
+    uintptr_t raddr = reinterpret_cast<uintptr_t>(remote_addr);
+    qps[peer]->atomic_add(raddr, rkey, value, wf_info);
     printf("gin_atomic_kernel: atomic posted, calling quiet\n");
     qps[peer]->quiet(wf_info);
     printf("gin_atomic_kernel: quiet done\n");
@@ -85,11 +88,15 @@ __global__ void gin_put_signal_kernel(rocshmem::QueuePair **qps,
                                        int peer) {
   if (threadIdx.x == 0 && blockIdx.x == 0) {
     rocshmem::ActiveWFInfo wf_info(peer, rocshmem::ThreadScope::thread);
+    static constexpr auto post_options = rocshmem::PostOpt{rocshmem::RingDB<false>};
     printf("gin_put_signal: peer=%d put dst=%p rkey=0x%x src=%p lkey=0x%x bytes=%zu\n",
            peer, dst, dst_rkey, src, src_lkey, nbytes);
-    qps[peer]->put_nbi(dst, dst_rkey, src, src_lkey, nbytes, wf_info, /*ring_db=*/false);
+    uintptr_t dst_raddr = reinterpret_cast<uintptr_t>(dst);
+    uintptr_t src_laddr = reinterpret_cast<uintptr_t>(src);
+    uintptr_t signal_raddr = reinterpret_cast<uintptr_t>(signal_addr);
+    qps[peer]->put_nbi(dst_raddr, dst_rkey, src_laddr, src_lkey, nbytes, wf_info, post_options);
     printf("gin_put_signal: atomic signal=%p rkey=0x%x\n", signal_addr, signal_rkey);
-    qps[peer]->atomic_add(signal_addr, signal_rkey, 1, wf_info);
+    qps[peer]->atomic_add(signal_raddr, signal_rkey, 1, wf_info);
     printf("gin_put_signal: calling quiet\n");
     qps[peer]->quiet(wf_info);
     printf("gin_put_signal: done\n");

projects/rccl/CHANGELOG.md

@@ -14,6 +14,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 * Multiple GIN proxy progress threads via `NCCL_GIN_PROXY_NTHREADS`.
 * Enabled hierarchical Copy Engine `ncclAllGather` and `ncclAlltoAll` 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. 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.
+* 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
 * **Breaking: `NCCL_GIN_TYPE` values for AMD backends are not compatible with 2.30.7.** NCCL 2.31 inserted EFA GDA at value 5, so rocSHMEM GDA moved 5→6 and Anvil SDMA moved 6→7. The IB proxy remains `2`. Jobs that still set `NCCL_GIN_TYPE=6` now select rocSHMEM GDA, not Anvil SDMA. See `src/gin/README.md`.

projects/rccl/cmake/DeviceLinker.cmake

@@ -287,6 +287,9 @@ foreach(DL_GPU_TARGET ${DL_GPU_TARGETS})
   if((ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN) AND TARGET rocshmem_static)
     add_dependencies(${_dev_target} rocshmem_static)
   endif()
+  if(ENABLE_ROCSHMEM_GIN AND TARGET copy_rocshmem_headers)
+    add_dependencies(${_dev_target} copy_rocshmem_headers)
+  endif()
   # Pass rocshmem device bitcode to per-kernel compiles so rocshmem device
   # symbols resolve during the per-arch device.elf link step.
   # ENABLE_ROCSHMEM: rocshmem_n_pes, alltoall_wg, etc.
@@ -1101,13 +1104,7 @@ if(GENERATE_SYM_KERNELS)
 
     add_custom_command(
       OUTPUT ${_qp_bc}
-      COMMAND ${_llvm_link}
-        ${_bc_dir}/queue_pair.bc
-        ${_bc_dir}/queue_pair_mlx5.bc
-        ${_bc_dir}/queue_pair_bnxt.bc
-        ${_bc_dir}/queue_pair_ionic.bc
-        ${_cm_bc}
-        -o ${_qp_raw}
+      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
@@ -1164,6 +1161,9 @@ add_custom_target(device_linker_build ALL
   DEPENDS ${COMMON_FAT_OBJ} ${ONERANK_FAT_OBJ} ${COLLECTIVES_FAT_OBJ} ${DIAG_P2P_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} ${RCCL_EP_FAT_OBJ} ${DEVICE_ELF_SYMLINKS}
 )
 add_dependencies(device_linker_build hipify_all copy_nccl_device_headers)
+if(ENABLE_ROCSHMEM_GIN AND TARGET copy_rocshmem_headers)
+  add_dependencies(device_linker_build copy_rocshmem_headers)
+endif()
 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.
@@ -1202,3 +1202,6 @@ set(DEVICE_LINKER_OBJECTS
 # ===========================================================================
 add_custom_target(device_ir DEPENDS ${ALL_IR_FILES})
 add_dependencies(device_ir hipify_all copy_nccl_device_headers)
+if(ENABLE_ROCSHMEM_GIN AND TARGET copy_rocshmem_headers)
+  add_dependencies(device_ir copy_rocshmem_headers)
+endif()

projects/rccl/cmake/ROCSHMEM.cmake

@@ -98,6 +98,10 @@ function(add_rocshmem_targets)
     # -----------------------------------------------------------------
     if(ROCSHMEM_INSTALL_DIR)
         list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+        # Set rocshmem_static_ROOT; variable read on enter, so must be set BEFORE find_package()
+        if(NOT rocshmem_static_ROOT)
+          set(rocshmem_static_ROOT "${ROCSHMEM_INSTALL_DIR}")
+        endif()
         find_package(rocshmem_static)
         if(rocshmem_static_FOUND)
             set(ROCSHMEM_INCLUDE_DIR "${ROCSHMEM_INCLUDE_DIR}" PARENT_SCOPE)
@@ -158,3 +162,23 @@ function(add_rocshmem_targets)
     set(ROCSHMEM_SOURCE_DIR "${ROCSHMEM_SOURCE_DIR}" PARENT_SCOPE)
 
 endfunction()
+
+# copy files in the list-valued variable ${FILE_LIST_VAR} from ${SRC_DIR} to ${DST_DIR}
+# appends the destination file location to the list-valued variable ${OUTPUT_LIST_VAR}
+function(copy_files FILE_LIST_VAR SRC_DIR DST_DIR OUTPUT_LIST_VAR)
+  foreach(file_name ${${FILE_LIST_VAR}})
+    set(src_file "${SRC_DIR}/${file_name}")
+    set(dst_file "${DST_DIR}/${file_name}")
+    get_filename_component(dst_file_dir "${dst_file}" DIRECTORY)
+    add_custom_command(
+      OUTPUT "${dst_file}"
+      COMMAND ${CMAKE_COMMAND} -E make_directory "${dst_file_dir}"
+      COMMAND ${CMAKE_COMMAND} -E copy_if_different "${src_file}" "${dst_file}"
+      DEPENDS "${src_file}"
+      COMMENT "Copying ${src_file} -> ${dst_file}"
+      VERBATIM
+    )
+    list(APPEND ${OUTPUT_LIST_VAR} "${dst_file}")
+  endforeach()
+  set(${OUTPUT_LIST_VAR} ${${OUTPUT_LIST_VAR}} PARENT_SCOPE)
+endfunction()

projects/rccl/src/CMakeLists.txt

@@ -265,7 +265,6 @@ set(SRC_FILES
   include/nccl_device/gin/proxy/gin_proxy_device_host_common.h
   include/nccl_device/gin/rocshmem_gda/gin_rocshmem_device_host_common_gda.h
   include/nccl_device/gin/rocshmem_gda/gin_rocshmem_gda.h
-  include/nccl_device/gin/rocshmem_gda/queue_pair_device.h
   include/nccl_device/gin_barrier.h
   include/nccl_device/impl/barrier__funcs.h
   include/nccl_device/impl/barrier__types.h
@@ -594,6 +593,13 @@ if(ENABLE_ROCSHMEM_GIN)
     src/algorithms/gin/sdma/gin_alltoall_sdma.cu
     src/algorithms/gin/sdma/gin_all_reduce_sdma.cu 
   )
+  list(APPEND SRC_FILES
+    src/include/nccl_device/gin/rocshmem_gda/log.hpp
+    src/include/nccl_device/gin/rocshmem_gda/util.hpp
+    src/include/nccl_device/gin/rocshmem_gda/containers/free_list.hpp
+    src/include/nccl_device/gin/rocshmem_gda/gda/ibv_core.hpp
+    src/include/nccl_device/gin/rocshmem_gda/rocshmem/rocshmem_common.hpp
+  )
 endif()
 
 if(USE_AMDSMI)
@@ -1055,6 +1061,90 @@ if (ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN)
     endif()
   endif()
 endif()
+
+if (ENABLE_ROCSHMEM_GIN)
+  # Ensure we see our copied headers before rocSHMEM's
+  target_include_directories(rccl BEFORE PRIVATE ${PROJECT_BINARY_DIR}/include/nccl_device/gin/rocshmem_gda)
+
+  # Need to add the GIN GDA directory to the public interface:
+  # internal rocSHMEM headers assume that the source directory is in the include search path
+  target_include_directories(rccl BEFORE INTERFACE
+    $<BUILD_INTERFACE:${PROJECT_BINARY_DIR}/include/nccl_device/gin/rocshmem_gda>
+    $<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/nccl_device/gin/rocshmem_gda>
+  )
+
+  # Add target for copying the rocSHMEM headers
+  add_custom_target(copy_rocshmem_headers)
+  add_dependencies(rccl copy_rocshmem_headers)
+
+  if (ROCSHMEM_INCLUDE_DIR)
+    # Copy rocSHMEM public include headers (primarily rocshmem_config.h) into RCCL
+    set(ROCSHMEM_INCLUDE_HEADERS rocshmem/rocshmem_config.h)
+    copy_files(
+      ROCSHMEM_INCLUDE_HEADERS
+      "${ROCSHMEM_INCLUDE_DIR}"
+      "${PROJECT_BINARY_DIR}/include/nccl_device/gin/rocshmem_gda"
+      GIN_GDA_ROCSHMEM_INCLUDE_HEADERS
+    )
+    add_custom_target(copy_rocshmem_include_headers DEPENDS ${GIN_GDA_ROCSHMEM_INCLUDE_HEADERS})
+    # Ensure rocSHMEM is fully built/installed before copying installed headers
+    if (TARGET rocshmem_ext)
+      add_dependencies(copy_rocshmem_include_headers rocshmem_ext)
+    endif()
+    add_dependencies(copy_rocshmem_headers copy_rocshmem_include_headers)
+  endif()
+
+  if (ROCSHMEM_SOURCE_DIR)
+    # Copy rocSHMEM QueuePair device headers into RCCL
+    set(ROCSHMEM_SOURCE_HEADERS
+      # Queue Pair core
+      gda/queue_pair/queue_pair_common.hpp
+      gda/queue_pair/queue_pair_option.hpp
+      gda/queue_pair/queue_pair_interface.hpp
+      gda/queue_pair/queue_pair_shmem.hpp
+      gda/queue_pair/queue_pair_device.hpp
+      gda/queue_pair_provider.hpp
+      # GDA_MUX (Multiplexing provider)
+      gda/queue_pair_mux.hpp
+      # GDA_IONIC (AMD Pensando IONIC RDMA provider)
+      gda/ionic/queue_pair_ionic.hpp
+      gda/ionic/provider_gda_ionic.hpp
+      gda/ionic/ionic_dv.h
+      gda/ionic/ionic_fw.h
+      # GDA_BNXT (Broadcom BNXT RDMA provider)
+      gda/bnxt/queue_pair_bnxt.hpp
+      gda/bnxt/provider_gda_bnxt.hpp
+      gda/bnxt/bnxt_re_dv.h
+      gda/bnxt/bnxt_re_hsi.h
+      # GDA_MLX5 (Mellanox MLX5 RDMA provider)
+      gda/mlx5/queue_pair_mlx5.hpp
+      gda/mlx5/provider_gda_mlx5.hpp
+      gda/mlx5/mlx5dv_core.hpp
+      gda/mlx5/mlx5_ifc_core.hpp
+      # GDA_QUEUEPAIR_MOCK (QueuePair mock object)
+      gda/queue_pair/queue_pair_mock.hpp
+      # rocSHMEM internals / miscellaneous
+      gda/endian.hpp
+      gda/gda_enums.hpp
+      bit.hpp
+      constmem.hpp
+      constants.hpp
+    )
+    copy_files(
+      ROCSHMEM_SOURCE_HEADERS
+      "${ROCSHMEM_SOURCE_DIR}/src"
+      "${PROJECT_BINARY_DIR}/include/nccl_device/gin/rocshmem_gda"
+      GIN_GDA_ROCSHMEM_SOURCE_HEADERS
+    )
+    add_custom_target(copy_rocshmem_source_headers DEPENDS ${GIN_GDA_ROCSHMEM_SOURCE_HEADERS})
+    add_dependencies(copy_rocshmem_headers copy_rocshmem_source_headers)
+  endif()
+
+  # Add rule to install ibvcore.h so that it can be found by ibv_core.hpp
+  rocm_install(FILES       ${HIPIFY_DIR}/src/include/ibvcore.h
+               DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/nccl_device/gin/rocshmem_gda)
+endif()
+
 if (ENABLE_ROCSHMEM)
   target_link_libraries(rccl PRIVATE ${IBVERBS})
 endif()