Skip to content
Merged
Show file tree
Hide file tree
Changes from 11 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 10 additions & 0 deletions projects/rccl/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -848,6 +848,16 @@ rocm_install(FILES ${PROJECT_BINARY_DIR}/include/nccl.h
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR})
rocm_install(DIRECTORY ${PROJECT_BINARY_DIR}/include/nccl_device/
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/nccl_device)
# Install SDMA device headers so consumers can compile GIN SDMA device templates
# without depending on ROCSHMEM_SOURCE_DIR.
if(ENABLE_ROCSHMEM_GIN AND ROCSHMEM_SOURCE_DIR)
Comment thread
abouteiller marked this conversation as resolved.
Resolved
Comment thread
abouteiller marked this conversation as resolved.
Resolved
rocm_install(FILES
Comment thread
abouteiller marked this conversation as resolved.
Resolved
${ROCSHMEM_SOURCE_DIR}/src/sdma/anvil_device.hpp
${ROCSHMEM_SOURCE_DIR}/src/sdma/sdma_opcodes.h
${ROCSHMEM_SOURCE_DIR}/src/sdma/sdma_pkt_struct.h
${ROCSHMEM_SOURCE_DIR}/src/sdma/sdma_pkt_struct_mi4.h
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/nccl_device/gin/anvil_sdma/sdma)
endif()
## Optional: install LLVM bitcode artifacts when EMIT_LLVM_IR is enabled.
## librccl_device.bc, nccl_device_wrapper.h, hipified nccl_device headers,
## and nccl.h (same as rccl.h; nccl_device #includes <nccl.h>) are produced
Expand Down
4 changes: 0 additions & 4 deletions projects/rccl/cmake/Findrocshmem_static.cmake
Comment thread
omor1 marked this conversation as resolved.
Resolved
Original file line number Diff line number Diff line change
Expand Up @@ -27,9 +27,5 @@ find_path(ROCSHMEM_INCLUDE_DIR
find_library(ROCSHMEM_LIBRARY
NAMES rocshmem
HINTS ${ROCSHMEM_INSTALL_DIR}/lib)

## -- todo --- what to do with verbs? add to handle args call below? -- ##
find_library(IBVERBS ibverbs)

find_package_handle_standard_args(rocshmem_static DEFAULT_MSG ROCSHMEM_INCLUDE_DIR ROCSHMEM_LIBRARY)
## mark_as_advanced(MSCCLPP_INCLUDE_DIRS MSCCLPP_NCCL_STATIC_LIB) add this for Rocshmem?
13 changes: 4 additions & 9 deletions projects/rccl/cmake/ROCSHMEM.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -29,13 +29,6 @@ set(ROCSHMEM_MONO_HASH "33d980d7ca1f0bf90cfe4ff9106310abcf47b550" CACHE STRING

function(add_rocshmem_targets)

# Common dependency: libibverbs is required for all rocSHMEM paths
find_library(_IBVERBS ibverbs)
if(NOT _IBVERBS)
message(FATAL_ERROR "libibverbs not found (install rdma-core/libibverbs-dev)")
endif()
set(IBVERBS ${_IBVERBS} PARENT_SCOPE)

# -----------------------------------------------------------------
# Auto-detect ROCSHMEM_SOURCE_DIR if not provided.
# Runs first so source headers are available regardless of whether
Expand Down Expand Up @@ -106,7 +99,7 @@ function(add_rocshmem_targets)
if(rocshmem_static_FOUND)
set(ROCSHMEM_INCLUDE_DIR "${ROCSHMEM_INCLUDE_DIR}" PARENT_SCOPE)
set(ROCSHMEM_LIBRARY "${ROCSHMEM_LIBRARY}" PARENT_SCOPE)
set(ROCSHMEM_SOURCE_DIR "${ROCSHMEM_SOURCE_DIR}" PARENT_SCOPE)
set(ROCSHMEM_SOURCE_DIR "${ROCSHMEM_SOURCE_DIR}" CACHE INTERNAL "rocSHMEM source directory")
return()
endif()
endif()
Expand Down Expand Up @@ -159,7 +152,9 @@ function(add_rocshmem_targets)
add_custom_target(rocshmem_static ALL DEPENDS rocshmem_ext)
endif()

set(ROCSHMEM_SOURCE_DIR "${ROCSHMEM_SOURCE_DIR}" PARENT_SCOPE)
# CACHE INTERNAL (not PARENT_SCOPE): SOURCE_DIR is consumed by the
# top-level CMakeLists.txt install rules, outside add_subdirectory(src).
set(ROCSHMEM_SOURCE_DIR "${ROCSHMEM_SOURCE_DIR}" CACHE INTERNAL "rocSHMEM source directory")

endfunction()

Expand Down
6 changes: 4 additions & 2 deletions projects/rccl/install.sh
Original file line number Diff line number Diff line change
Expand Up @@ -328,11 +328,13 @@ fi
# rocSHMEM worktree setup (must run before cd-ing into the build directory)
# #################################################
rocshmem_source_dir=""
rocshmem_is_monorepo=false
if [[ "${build_rocshmem_support}" == true || "${build_rocshmem_gin}" == true ]] && [[ -z "${ROCSHMEM_INSTALL_DIR}" ]]; then
# Prefer mono-repo layout (projects/rocshmem alongside projects/rccl)
mono_root=$(git rev-parse --show-toplevel 2>/dev/null)
if [[ -n "$mono_root" ]] && [[ -f "$mono_root/projects/rocshmem/CMakeLists.txt" ]]; then
rocshmem_source_dir="$mono_root/projects/rocshmem"
rocshmem_is_monorepo=true
echo "=== Using rocSHMEM from mono-repo: ${rocshmem_source_dir} ==="
else
setup_rocshmem_worktree
Expand Down Expand Up @@ -511,14 +513,14 @@ if [[ "${build_rocshmem_support}" == true ]]; then
cmake_common_options="${cmake_common_options} -DENABLE_ROCSHMEM=ON"
if [[ -n "${ROCSHMEM_INSTALL_DIR}" ]]; then
cmake_common_options="${cmake_common_options} -DROCSHMEM_INSTALL_DIR=${ROCSHMEM_INSTALL_DIR}"
elif [[ -n "${rocshmem_source_dir}" ]]; then
elif [[ -n "${rocshmem_source_dir}" ]] && [[ "${rocshmem_is_monorepo}" == false ]]; then
cmake_common_options="${cmake_common_options} -DROCSHMEM_SOURCE_DIR=${rocshmem_source_dir}"
fi
elif [[ "${build_rocshmem_gin}" == true ]]; then
cmake_common_options="${cmake_common_options} -DENABLE_ROCSHMEM=OFF -DENABLE_ROCSHMEM_GIN=ON"
if [[ -n "${ROCSHMEM_INSTALL_DIR}" ]]; then
cmake_common_options="${cmake_common_options} -DROCSHMEM_INSTALL_DIR=${ROCSHMEM_INSTALL_DIR}"
elif [[ -n "${rocshmem_source_dir}" ]]; then
elif [[ -n "${rocshmem_source_dir}" ]] && [[ "${rocshmem_is_monorepo}" == false ]]; then
cmake_common_options="${cmake_common_options} -DROCSHMEM_SOURCE_DIR=${rocshmem_source_dir}"
fi
else
Expand Down
31 changes: 21 additions & 10 deletions projects/rccl/src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -214,6 +214,8 @@ set(SRC_FILES
include/gin/gin_host_proxy.h
include/gin/gin_host_rocshmem_gda.h
include/gin/gin_rocshmem_gda_factory.h
include/gin/gin_host_anvil_sdma.h
include/gin/gin_anvil_sdma_factory.h
include/git_version.h
include/graph.h
include/group.h
Expand Down Expand Up @@ -589,17 +591,22 @@ if(ENABLE_ROCSHMEM_GIN)
src/gin/gin_plugin_rocshmem_gda.cc
src/gin/gin_rocshmem_gda_factory.cc
src/gin/gin_plugin_anvil_sdma.cc
src/gin/gin_anvil_sdma_factory.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
src/algorithms/gin/sdma/gin_all_reduce_sdma.cu
)
# Rocshmem GDA shim headers: in SRC_FILES so hipify copies them to the
# hipify output tree (hipify is a no-op on them — no CUDA types).
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
)
# anvil.cpp is compiled from rocshmem source via target_sources after
# add_library (bypasses SRC_FILES/hipify). See block after add_library.
endif()

if(USE_AMDSMI)
Expand Down Expand Up @@ -1059,6 +1066,16 @@ if (ENABLE_ROCSHMEM OR ENABLE_ROCSHMEM_GIN)
else()
target_include_directories(rccl PRIVATE ${ROCSHMEM_SOURCE_DIR}/build/include ${ROCSHMEM_SOURCE_DIR}/build/include/rocshmem)
endif()

# Compile anvil.cpp from rocshmem source tree into librccl.so (host-only).
# src/gin/log.hpp shim shadows rocshmem's log.hpp via include path ordering,
Comment thread
abouteiller marked this conversation as resolved.
Resolved
# redirecting LOG_* macros to RCCL native logging (WARN/INFO/TRACE).
if(ENABLE_ROCSHMEM_GIN)
target_sources(rccl PRIVATE "${ROCSHMEM_SOURCE_DIR}/src/sdma/anvil.cpp")
set_source_files_properties("${ROCSHMEM_SOURCE_DIR}/src/sdma/anvil.cpp" PROPERTIES
INCLUDE_DIRECTORIES "${CMAKE_CURRENT_SOURCE_DIR}/gin;${ROCSHMEM_SOURCE_DIR}/src/sdma;${ROCSHMEM_SOURCE_DIR}/src;${ROCSHMEM_SOURCE_DIR}/include;${ROCM_PATH}/include")
Comment thread
abouteiller marked this conversation as resolved.
Resolved
endif()

endif()
endif()

Expand Down Expand Up @@ -1492,7 +1509,6 @@ if(ENABLE_ROCSHMEM)
# of the final librccl.so link.
target_link_libraries(rccl PRIVATE ${ROCSHMEM_LIBRARY})
target_link_libraries(rccl PRIVATE ${HSAKMT_LIBRARY})
target_link_libraries(rccl PRIVATE ${IBVERBS})
set(_rocshmem_offload_arch_flags "")
foreach(_gpu_raw ${GPU_TARGETS})
string(REGEX REPLACE ":.*" "" _gpu "${_gpu_raw}")
Expand All @@ -1505,15 +1521,10 @@ elseif(ENABLE_ROCSHMEM_GIN)
# With the device linker (default), per-arch .bc files are injected directly
# into the per-arch device link step via DeviceLinker.cmake.
#
# 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.
# Seed rocshmem host symbols needed by GDA into librccl.so. SDMA factory
# symbols are compiled in directly (gin_anvil_sdma_factory.cc) and no longer
# need pulling from librocshmem.a.
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})
Expand Down
Original file line number Diff line number Diff line change
@@ -1,20 +1,21 @@
/******************************************************************************
/*************************************************************************
* Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
*
* SPDX-License-Identifier: MIT
*****************************************************************************/
* See LICENSE.txt for license information
************************************************************************/

#include <gin_anvil/sdma_factory.h>
#ifdef ENABLE_ROCSHMEM_GIN

#include "anvil.hpp"
#include "util.hpp"
#include "gin/gin_anvil_sdma_factory.h"

#include "sdma/anvil.hpp"
#include <hip/hip_runtime.h>
#include <cstdlib>
#include <vector>
#include <cstring>
#include <algorithm>

#include "log.hpp"
#include "debug.h"

struct gin_anvil_sdma_opaque {
int nRanks;
Expand All @@ -34,7 +35,7 @@ extern "C" int gin_anvil_sdma_probe(void) {

static int checkHip(hipError_t e, const char* what) {
if (e != hipSuccess) {
LOG_ERROR("gin_anvil_sdma: %s: %s", what, hipGetErrorString(e));
WARN("GIN anvil-sdma: %s: %s", what, hipGetErrorString(e));
return -1;
}
return 0;
Expand All @@ -61,21 +62,21 @@ extern "C" int gin_anvil_sdma_create(int nRanks, int myRank, int my_device_id,
std::vector<int> devs(static_cast<size_t>(nRanks), -1);
devs[static_cast<size_t>(myRank)] = my_device_id;
if (allgather(allgather_ctx, devs.data(), sizeof(int)) != 0) {
Comment thread
abouteiller marked this conversation as resolved.
Resolved
LOG_ERROR("gin_anvil_sdma: allgather(device ids) failed");
WARN("GIN anvil-sdma: allgather(device ids) failed");
return -1;
}

for (int i = 0; i < nRanks; ++i) {
if (devs[static_cast<size_t>(i)] < 0) {
Comment thread
abouteiller marked this conversation as resolved.
Resolved
LOG_ERROR("gin_anvil_sdma: invalid device id for rank %d", i);
WARN("GIN anvil-sdma: invalid device id for rank %d", i);
return -1;
}
}

const int myDev = devs[static_cast<size_t>(myRank)];

if (!sdma_anvil::initEndpoint()) {
LOG_ERROR("gin_anvil_sdma: Anvil SDMA init failed");
WARN("GIN anvil-sdma: Anvil SDMA init failed");
return -1;
}

Expand All @@ -85,12 +86,12 @@ extern "C" int gin_anvil_sdma_create(int nRanks, int myRank, int my_device_id,
if (sdma_anvil::anvil.getSdmaQueue(myDev, remoteDev, 0) != nullptr) continue;
if (myDev != remoteDev) sdma_anvil::EnablePeerAccess(myDev, remoteDev);
if (!sdma_anvil::anvil.connect(myDev, remoteDev, numChannels)) {
LOG_ERROR("gin_anvil_sdma: connect(%d -> %d) failed", myDev, remoteDev);
WARN("GIN anvil-sdma: connect(%d -> %d) failed", myDev, remoteDev);
return -1;
}
}
} catch (const std::exception& e) {
LOG_ERROR("gin_anvil_sdma: SDMA connect failed: %s", e.what());
WARN("GIN anvil-sdma: SDMA connect failed: %s", e.what());
return -1;
}

Expand All @@ -108,7 +109,7 @@ extern "C" int gin_anvil_sdma_create(int nRanks, int myRank, int my_device_id,
}
}
if (validHandles == 0) {
LOG_ERROR("gin_anvil_sdma: no SDMA queue handles for device %d", myDev);
WARN("GIN anvil-sdma: no SDMA queue handles for device %d", myDev);
return -1;
}

Expand All @@ -118,19 +119,19 @@ extern "C" int gin_anvil_sdma_create(int nRanks, int myRank, int my_device_id,
if (checkHip(hipMemcpy(dev_row, host_handles.data(), static_cast<size_t>(total) * sizeof(void*),
hipMemcpyHostToDevice),
"hipMemcpy handles") != 0) {
CHECK_HIP(hipFree(dev_row));
checkHip(hipFree(dev_row), "hipFree handles (cleanup)");
return -1;
}

uint64_t* dirty = nullptr;
if (checkHip(hipExtMallocWithFlags((void**)&dirty, sizeof(uint64_t), hipDeviceMallocFinegrained),
"hipExtMallocWithFlags sdmaDirty") != 0) {
CHECK_HIP(hipFree(dev_row));
checkHip(hipFree(dev_row), "hipFree handles (cleanup)");
return -1;
}
if (checkHip(hipMemset(dirty, 0, sizeof(uint64_t)), "hipMemset sdmaDirty") != 0) {
CHECK_HIP(hipFree(dev_row));
CHECK_HIP(hipFree(dirty));
checkHip(hipFree(dev_row), "hipFree handles (cleanup)");
checkHip(hipFree(dirty), "hipFree dirty (cleanup)");
return -1;
}

Expand All @@ -152,8 +153,8 @@ extern "C" int gin_anvil_sdma_create(int nRanks, int myRank, int my_device_id,
extern "C" void gin_anvil_sdma_destroy(gin_anvil_sdma_handle_t handle) {
if (!handle) return;
auto* impl = handle;
if (impl->deviceHandles_d) CHECK_HIP(hipFree(impl->deviceHandles_d));
if (impl->sdmaDirty_d) CHECK_HIP(hipFree(impl->sdmaDirty_d));
if (impl->deviceHandles_d) checkHip(hipFree(impl->deviceHandles_d), "hipFree handles (destroy)");
if (impl->sdmaDirty_d) checkHip(hipFree(impl->sdmaDirty_d), "hipFree dirty (destroy)");
delete impl;
}

Expand All @@ -168,3 +169,5 @@ extern "C" int gin_anvil_sdma_get_num_channels(gin_anvil_sdma_handle_t handle) {
extern "C" int gin_anvil_sdma_get_channel_stride(gin_anvil_sdma_handle_t handle) {
return handle ? handle->sdmaChannelStride : 0;
}

#endif // ENABLE_ROCSHMEM_GIN
2 changes: 1 addition & 1 deletion projects/rccl/src/gin/gin_plugin_anvil_sdma.cc
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
#include "bootstrap.h"
#include "nccl_device/gin/anvil_sdma/gin_anvil_sdma_device_host_common.h"
#include "nccl_device/gin/anvil_sdma/gin_anvil_ipc_table.h"
#include <gin_anvil/sdma_factory.h>
#include "gin/gin_anvil_sdma_factory.h"
#include <hip/hip_runtime.h>
#include <cstdint>
#include <cstdlib>
Expand Down
2 changes: 1 addition & 1 deletion projects/rccl/src/gin/gin_rocshmem_gda_factory.cc
Original file line number Diff line number Diff line change
Expand Up @@ -39,9 +39,9 @@
*/

#include "gin/gin_rocshmem_gda_factory.h"
#include <gda/gda_enums.hpp>
#include <gda/queue_pair_provider.hpp>
#include <gda/ibv_wrapper.hpp>
#include <gda/gda_enums.hpp>
Comment thread
abouteiller marked this conversation as resolved.
Resolved
#include <gda/debug_gda.hpp>
#include <envvar.hpp>
#include <gda/topology.hpp>
Expand Down
49 changes: 49 additions & 0 deletions projects/rccl/src/gin/log.hpp
Comment thread
omor1 marked this conversation as resolved.
Resolved
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
/*************************************************************************
* Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
*
* See LICENSE.txt for license information
************************************************************************/

// Shim: shadows rocshmem's log.hpp when compiling anvil.cpp into librccl.so.
// Maps rocshmem LOG_* macros to RCCL's native logging infrastructure.

#ifndef RCCL_GIN_LOG_SHIM_HPP
#define RCCL_GIN_LOG_SHIM_HPP

// Host-side: when building librccl.so, map to RCCL native logging.
// When installed (consumed by rccl-tests), host macros are not needed
// — only device-side stubs below are used by anvil_device.hpp.
#if __has_include("debug.h")
Comment thread
abouteiller marked this conversation as resolved.
Resolved
#include "debug.h"
#include <cstdlib>
#define LOG_ERROR(fmt, ...) WARN("GIN: " fmt, ##__VA_ARGS__)
#define LOG_ERROR_EXIT(fmt, ...) do { WARN("GIN FATAL: " fmt, ##__VA_ARGS__); exit(EXIT_FAILURE); } while(0)
#define LOG_ERROR_ABORT(fmt, ...) do { WARN("GIN ABORT: " fmt, ##__VA_ARGS__); abort(); } while(0)
#define LOG_WARN(fmt, ...) INFO(NCCL_INIT, "GIN: " fmt, ##__VA_ARGS__)
#define LOG_INFO(fmt, ...) INFO(NCCL_INIT, "GIN: " fmt, ##__VA_ARGS__)
#ifdef ENABLE_TRACE
#define LOG_TRACE(fmt, ...) TRACE(NCCL_INIT, "GIN: " fmt, ##__VA_ARGS__)
#else
#define LOG_TRACE(...) do {} while(0)
#endif
#else
// Installed context: no RCCL internals available.
#include <cstdio>
Comment thread
abouteiller marked this conversation as resolved.
Resolved
#include <cstdlib>
#define LOG_ERROR(fmt, ...) fprintf(stderr, "GIN: " fmt "\n", ##__VA_ARGS__)
#define LOG_ERROR_EXIT(fmt, ...) do { fprintf(stderr, "GIN FATAL: " fmt "\n", ##__VA_ARGS__); exit(EXIT_FAILURE); } while(0)
#define LOG_ERROR_ABORT(fmt, ...) do { fprintf(stderr, "GIN ABORT: " fmt "\n", ##__VA_ARGS__); abort(); } while(0)
#define LOG_WARN(fmt, ...) fprintf(stderr, "GIN: " fmt "\n", ##__VA_ARGS__)
#define LOG_INFO(fmt, ...) ((void)0)
Comment thread
abouteiller marked this conversation as resolved.
Resolved
#define LOG_TRACE(...) do {} while(0)
#endif

// Device-side logging stubs (anvil_device.hpp uses LOGD_ERROR_ABORT for
// unsupported-architecture guards — map to __builtin_trap on device).
#define LOGD_ERROR_ABORT(fmt, ...) __builtin_trap()
Comment thread
abouteiller marked this conversation as resolved.
Resolved
#define LOGD_ERROR(fmt, ...) ((void)0)
Comment thread
omor1 marked this conversation as resolved.
Resolved
#define LOGD_WARN(fmt, ...) ((void)0)
#define LOGD_INFO(fmt, ...) ((void)0)
#define LOGD_TRACE(fmt, ...) ((void)0)

#endif // RCCL_GIN_LOG_SHIM_HPP
Loading
Loading