diff --git a/CMakeLists.txt b/CMakeLists.txt index af9197e5..8a4b7f65 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -199,9 +199,7 @@ include(CMakePushCheckState) # Build options #================================================================================================== option(BUILD_LOCAL_GPU_TARGET_ONLY "Build only for GPUs detected on this machine" OFF) -option(ENABLE_NIC_EXEC "Enable RDMA NIC Executor in TransferBench" OFF) option(ENABLE_MPI_COMM "Enable MPI Communicator support" OFF) -option(ENABLE_DMA_BUF "Enable DMA-BUF support for GPU Direct RDMA" OFF) option(ENABLE_AMD_SMI "Enable AMD-SMI pod membership queries" OFF) option(ENABLE_POD_COMM "Enable pod communication" OFF) option(BUILD_RELOCATABLE_PACKAGE "Build with RVS-style relocatable RPATH and amdrocm-transferbench package naming" OFF) @@ -313,71 +311,13 @@ else() message(FATAL_ERROR "HSA library or headers not found under ${ROCM_PATH}; TransferBench requires libhsa-runtime64") endif() -## Check for infiniband verbs support -if(DEFINED ENV{DISABLE_NIC_EXEC} AND "$ENV{DISABLE_NIC_EXEC}" STREQUAL "1") - message(STATUS "Disabling NIC Executor support as env. flag DISABLE_NIC_EXEC was enabled") -elseif(NOT ENABLE_NIC_EXEC) - message(STATUS "For CMake builds, NIC Executor support requires explicit opt-in by setting CMake flag -DENABLE_NIC_EXEC=ON") - message(STATUS "- Disabling NIC Executor support") -else() - message(STATUS "Attempting to build with NIC executor support") - - find_library(IBVERBS_LIBRARY ibverbs) - find_path(IBVERBS_INCLUDE_DIR infiniband/verbs.h) - if(IBVERBS_LIBRARY AND IBVERBS_INCLUDE_DIR) - add_library(ibverbs SHARED IMPORTED) - set_target_properties(ibverbs PROPERTIES INTERFACE_INCLUDE_DIRECTORIES "${IBVERBS_INCLUDE_DIR}" IMPORTED_LOCATION "${IBVERBS_LIBRARY}" INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${IBVERBS_INCLUDE_DIR}") - set(IBVERBS_FOUND 1) - message(STATUS "- Building with NIC executor support. Can set DISABLE_NIC_EXEC=1 to disable") - else() - if(NOT IBVERBS_LIBRARY) - message(WARNING "- IBVerbs library not found") - elseif(NOT IBVERBS_INCLUDE_DIR) - message(WARNING "- infiniband/verbs.h not found") - endif() - message(WARNING "- Building without NIC executor support. To use the TransferBench RDMA executor, check if your system has NICs, the NIC drivers are installed, and libibverbs-dev is installed") - endif() -endif() - -## Check for DMA-BUF support (requires IBVERBS) -if(IBVERBS_FOUND) - if(DEFINED ENV{DISABLE_DMA_BUF} AND "$ENV{DISABLE_DMA_BUF}" STREQUAL "1") - message(STATUS "Disabling DMA-BUF support as env. flag DISABLE_DMA_BUF was enabled") - elseif(NOT ENABLE_DMA_BUF) - message(STATUS "For CMake builds, DMA-BUF support requires explicit opt-in by setting CMake flags -DENABLE_DMA_BUF=ON") - message(STATUS "- Disabling DMA-BUF support") - else() - message(STATUS "Attempting to build with DMA-BUF support") - - # Check for ibv_reg_dmabuf_mr - cmake_push_check_state() - set(CMAKE_REQUIRED_INCLUDES ${IBVERBS_INCLUDE_DIR}) - set(CMAKE_REQUIRED_LIBRARIES ${IBVERBS_LIBRARY}) - check_symbol_exists(ibv_reg_dmabuf_mr "infiniband/verbs.h" HAVE_IBV_DMABUF) - cmake_pop_check_state() - - # Check for hsa_amd_portable_export_dmabuf - cmake_push_check_state() - set(CMAKE_REQUIRED_INCLUDES ${HSA_INCLUDE_DIR}) - set(CMAKE_REQUIRED_LIBRARIES ${HSA_LIBRARY}) - check_symbol_exists(hsa_amd_portable_export_dmabuf "hsa/hsa_ext_amd.h" HAVE_ROCM_DMABUF) - cmake_pop_check_state() - - # Enable DMA-BUF only if both APIs are available - if(HAVE_IBV_DMABUF AND HAVE_ROCM_DMABUF) - set(DMABUF_SUPPORT_FOUND 1) - message(STATUS "- Building with DMA-BUF support") - else() - if(NOT HAVE_IBV_DMABUF AND NOT HAVE_ROCM_DMABUF) - message(WARNING "- Building without DMA-BUF support: missing both ibv_reg_dmabuf_mr and ROCm DMA-BUF export") - elseif(NOT HAVE_IBV_DMABUF) - message(WARNING "- Building without DMA-BUF support: missing ibv_reg_dmabuf_mr") - else() - message(WARNING "- Building without DMA-BUF support: missing ROCm DMA-BUF export") - endif() - endif() - endif() -endif() +## NIC / RDMA executor +## +## libibverbs is loaded dynamically at runtime via dlopen/dlsym +## (see third-party/ibverbs/IbvDynLoad.hpp). The build neither links against +## -libverbs nor requires libibverbs-dev to be installed on the build host. +## Only libdl needs to be linked so that dlopen/dlsym resolve. +message(STATUS "NIC executor: libibverbs is loaded dynamically at runtime (no -libverbs link, no build-host dependency)") ## Check for MPI support set(MPI_PATH "" CACHE PATH "Path to MPI installation (takes priority over system MPI)") @@ -510,13 +450,11 @@ add_executable(TransferBench src/client/Client.cpp) target_include_directories(TransferBench PRIVATE src/header src/client - src/client/Presets) + src/client/Presets + third-party/ibverbs) -if(IBVERBS_FOUND) - target_include_directories(TransferBench PRIVATE ${IBVERBS_INCLUDE_DIR}) - target_link_libraries(TransferBench PRIVATE ${IBVERBS_LIBRARY}) - target_compile_definitions(TransferBench PRIVATE NIC_EXEC_ENABLED) -endif() +# libdl supplies dlopen/dlsym used by third-party/ibverbs/IbvDynLoad.hpp. +target_link_libraries(TransferBench PRIVATE ${CMAKE_DL_LIBS}) if(MPI_COMM_FOUND) if(TARGET MPI::MPI_CXX) target_link_libraries(TransferBench PRIVATE MPI::MPI_CXX) @@ -526,9 +464,6 @@ if(MPI_COMM_FOUND) endif() target_compile_definitions(TransferBench PRIVATE MPI_COMM_ENABLED) endif() -if(DMABUF_SUPPORT_FOUND) - target_compile_definitions(TransferBench PRIVATE HAVE_DMABUF_SUPPORT) -endif() if(AMD_SMI_FOUND) target_include_directories(TransferBench PRIVATE ${AMD_SMI_INCLUDE_DIR}) target_link_libraries(TransferBench PRIVATE ${AMD_SMI_LIBRARY}) diff --git a/Makefile b/Makefile index db40ed20..1e52f026 100644 --- a/Makefile +++ b/Makefile @@ -6,14 +6,20 @@ ROCM_PATH ?= /opt/rocm CUDA_PATH ?= /usr/local/cuda MPI_PATH ?= /usr/local/openmpi +# pip ROCm wheels ship bin/amdclang++ but often omit bin/amdllvm (which that stub execs). +# Default to llvm/bin when bin/amdllvm is absent so HIP builds work without extra flags. +ifeq ("$(shell test -e $(ROCM_PATH)/bin/amdllvm && echo found)", "found") HIPCC ?= $(ROCM_PATH)/bin/amdclang++ +else ifeq ("$(shell test -e $(ROCM_PATH)/llvm/bin/amdclang++ && echo found)", "found") +HIPCC ?= $(ROCM_PATH)/llvm/bin/amdclang++ +else +HIPCC ?= $(ROCM_PATH)/bin/amdclang++ +endif NVCC ?= $(CUDA_PATH)/bin/nvcc DEBUG ?= 0 # Optional features (set to 0 to disable, 1 to enable) -# DISABLE_NIC_EXEC: Disable RDMA/NIC executor support (default: 0) # DISABLE_MPI_COMM: Disable MPI communicator support (default: 0) -# DISABLE_DMA_BUF: Disable DMA-BUF support for GPU Direct RDMA (default: 1) # DISABLE_AMD_SMI: Disable AMD-SMI pod membership checking support (default: 0) # DISABLE_NVML: Disable NVML pod membership detection for CUDA builds (default: 0) # DISABLE_POD_COMM: Disable pod communication support (default: 0) @@ -83,55 +89,13 @@ ifeq ($(filter clean,$(MAKECMDGOALS)),) else COMMON_FLAGS += -O0 -g -ggdb3 endif - COMMON_FLAGS += -I./src/header -I./src/client -I./src/client/Presets + COMMON_FLAGS += -I./src/header -I./src/client -I./src/client/Presets -I./third-party/ibverbs - LDFLAGS += -lpthread - - NIC_ENABLED = 0 - # Compile RDMA executor if - # 1) DISABLE_NIC_EXEC is not set to 1 - # 2) IBVerbs is found in the Dynamic Linker cache - # 3) infiniband/verbs.h is found in the default include path - DISABLE_NIC_EXEC ?= 0 - ifneq ($(DISABLE_NIC_EXEC),1) - $(info Attempting to build with NIC executor support) - ifeq ("$(shell ldconfig -p | grep -c ibverbs)", "0") - $(info - ibverbs library not found) - else ifeq ("$(shell echo '#include ' | $(CXX) -E - 2>/dev/null | grep -c 'infiniband/verbs.h')", "0") - $(info - infiniband/verbs.h not found) - else - COMMON_FLAGS += -DNIC_EXEC_ENABLED - LDFLAGS += -libverbs - NIC_ENABLED = 1 - - # Disable DMA-BUF support by default (set DISABLE_DMA_BUF=0 to enable) - DISABLE_DMA_BUF ?= 1 - ifeq ($(DISABLE_DMA_BUF), 0) - # Check for both ibv_reg_dmabuf_mr and ROCm DMA-BUF export support - HAVE_IBV_DMABUF := $(shell echo '#include ' | $(CXX) -E - 2>/dev/null | grep -c 'ibv_reg_dmabuf_mr') - HAVE_ROCM_DMABUF := $(shell echo '#include ' | $(CXX) -I$(ROCM_PATH)/include -E - 2>/dev/null | grep -c 'hsa_amd_portable_export_dmabuf') - - ifeq ($(HAVE_IBV_DMABUF):$(HAVE_ROCM_DMABUF), 0:0) - $(info Building without DMA-BUF support: missing both ibv_reg_dmabuf_mr and ROCm DMA-BUF export) - else ifeq ($(HAVE_IBV_DMABUF), 0) - $(info Building without DMA-BUF support: missing ibv_reg_dmabuf_mr) - else ifeq ($(HAVE_ROCM_DMABUF), 0) - $(info Building without DMA-BUF support: missing ROCm DMA-BUF export) - else - COMMON_FLAGS += -DHAVE_DMABUF_SUPPORT - $(info Building with DMA-BUF support) - endif - else - $(info Building with DMA-BUF support disabled (DISABLE_DMA_BUF=1)) - endif - endif - ifeq ($(NIC_ENABLED), 0) - $(info - Building without NIC executor support) - $(info - To use the TransferBench RDMA executor, check if your system has NICs, the NIC drivers are installed, and libibverbs-dev is installed) - else - $(info - Building with NIC executor support. Can set DISABLE_NIC_EXEC=1 to disable) - endif - endif + # libibverbs is loaded dynamically at runtime via dlopen/dlsym (see + # third-party/ibverbs/IbvDynLoad.hpp), so the build never links against -libverbs + # and does not require libibverbs-dev to be installed. We only need -ldl so + # the dynamic loader API is resolvable. + LDFLAGS += -lpthread -ldl MPI_ENABLED = 0 # Compile with MPI communicator support if diff --git a/build_packages_local.sh b/build_packages_local.sh index 91ace56b..e2c073d1 100755 --- a/build_packages_local.sh +++ b/build_packages_local.sh @@ -229,9 +229,7 @@ CMAKE_ARGS=( -DCMAKE_VERBOSE_MAKEFILE=ON -DBUILD_RELOCATABLE_PACKAGE=ON -DBUILD_LOCAL_GPU_TARGET_ONLY=OFF - -DENABLE_NIC_EXEC=OFF -DENABLE_MPI_COMM=OFF - -DDISABLE_DMABUF=OFF -DGPU_TARGETS="${GPU_TARGETS}" -DTRANSFERBENCH_PACKAGE_RELEASE="${PKG_RELEASE}" ) diff --git a/src/client/Client.cpp b/src/client/Client.cpp index 95a0f9a5..d8496eb7 100644 --- a/src/client/Client.cpp +++ b/src/client/Client.cpp @@ -226,9 +226,7 @@ int main(int argc, char **argv) void DisplayVersion() { bool nicSupport = false, mpiSupport = false, podSupport = false; -#if NIC_EXEC_ENABLED - nicSupport = true; -#endif + nicSupport = IsIbvSymbolsReady(); #if MPI_COMM_ENABLED mpiSupport = true; #endif diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 34940b69..f69fc9b5 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -124,6 +124,13 @@ class EnvVars int nicTrafficClass; // DSCP/traffic class byte for RoCE GRH int roceVersion; // RoCE version number + // TDM (async tensor) options + int tdmBlockSize; // Size of each threadblock for TDM (async tensor) kernels (must be multiple of 32) + int tdmMaxLdsBytes; // Max LDS (shared memory) bytes per workgroup for TDM kernels (INT_MAX = use device max) + int tdmBlockOrder; // Worker ranking across blocks for TDM kernels (0=sequential, 1=interleaved) + int tdmTemporal; // TDM non-temporal cache hint (0=none, 1=loads, 2=stores, 3=both) + int tdmWordSize; // TDM tensor descriptor data-size encoding (advanced; default 2) + // Developer features int gpuMaxHwQueues; // Tracks GPU_MAX_HW_QUEUES environment variable @@ -173,6 +180,11 @@ class EnvVars showBorders = GetEnvVar("SHOW_BORDERS" , 1); showIterations = GetEnvVar("SHOW_ITERATIONS" , 0); showPercentiles = GetEnvVarArray("SHOW_PERCENTILES", {}); + tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); + tdmMaxLdsBytes = GetEnvVar("TDM_MAX_LDS_BYTES" , INT_MAX); + tdmBlockOrder = GetEnvVar("TDM_BLOCK_ORDER" , 0); + tdmTemporal = GetEnvVar("TDM_TEMPORAL" , 0); + tdmWordSize = GetEnvVar("TDM_WORD_SIZE" , 2); useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); @@ -369,32 +381,35 @@ class EnvVars printf(" GFX_WORD_SIZE - GFX kernel packed data size (4=DWORDx4, 2=DWORDx2, 1=DWORDx1)\n"); printf(" GPU_MAX_HW_QUEUES - Max hardware queues per GPU device (default = 4)\n"); printf(" HIDE_ENV - Hide environment variable value listing\n"); -#if NIC_EXEC_ENABLED - printf(" IB_GID_INDEX - Required for RoCE NICs (default=-1/auto)\n"); - printf(" IB_PORT_NUMBER - RDMA port count for RDMA NIC (default=1)\n"); - printf(" IP_ADDRESS_FAMILY - IP address family (4=v4, 6=v6, default=v4)\n"); -#endif + if (IsIbvSymbolsReady()) { + printf(" IB_GID_INDEX - Required for RoCE NICs (default=-1/auto)\n"); + printf(" IB_PORT_NUMBER - RDMA port count for RDMA NIC (default=1)\n"); + printf(" IP_ADDRESS_FAMILY - IP address family (4=v4, 6=v6, default=v4)\n"); + printf(" NIC_CHUNK_BYTES - Number of bytes to send at a time using NIC (default = 1GB)\n"); + printf(" NIC_CQ_POLL_BATCH - Number of CQ entries to poll per ibv_poll_cq call (default = 4)\n"); + printf(" NIC_RELAX_ORDER - Set to non-zero to use relaxed ordering\n"); + printf(" NIC_SERVICE_LEVEL - IB service level (sl) for InfiniBand QPs (default=0)\n"); + printf(" NIC_TRAFFIC_CLASS - DSCP/traffic class byte for RoCE GRH (default=0)\n"); + printf(" ROCE_VERSION - RoCE version (default=2)\n"); + } printf(" MIN_VAR_SUBEXEC - Minimum # of subexecutors to use for variable subExec Transfers\n"); printf(" MAX_VAR_SUBEXEC - Maximum # of subexecutors to use for variable subExec Transfers (0 for device limits)\n"); -#if NIC_EXEC_ENABLED - printf(" NIC_CHUNK_BYTES - Number of bytes to send at a time using NIC (default = 1GB)\n"); - printf(" NIC_CQ_POLL_BATCH - Number of CQ entries to poll per ibv_poll_cq call (default = 4)\n"); - printf(" NIC_RELAX_ORDER - Set to non-zero to use relaxed ordering\n"); - printf(" NIC_SERVICE_LEVEL - IB service level (sl) for InfiniBand QPs (default=0)\n"); - printf(" NIC_TRAFFIC_CLASS - DSCP/traffic class byte for RoCE GRH (default=0)\n"); -#endif printf(" NUM_ITERATIONS - # of timed iterations per test. If negative, run for this many seconds instead\n"); printf(" NUM_SUBITERATIONS - # of sub-iterations to run per iteration. Must be non-negative\n"); printf(" NUM_WARMUPS - # of untimed warmup iterations per test\n"); printf(" OUTPUT_TO_CSV - Outputs to CSV format if set\n"); -#if NIC_EXEC_ENABLED - printf(" ROCE_VERSION - RoCE version (default=2)\n"); -#endif printf(" SAMPLING_FACTOR - Add this many samples (when possible) between powers of 2 when auto-generating data sizes\n"); printf(" SHOW_BORDERS - Show ASCII box-drawing characters in tables\n"); printf(" SHOW_ITERATIONS - Show per-iteration timing info\n"); printf(" SHOW_PERCENTILES - Comma-separated percentiles iteration duration\n"); - printf(" USE_HIP_EVENTS - Use HIP events for GFX executor timing\n"); + printf(" TDM_BLOCK_ORDER - TDM worker ranking across blocks. 0=sequential, 1=interleaved\n"); + printf(" TDM_BLOCK_SIZE - # of threads per threadblock for TDM (async tensor) kernels (Must be multiple of 32)\n"); + printf(" TDM_MAX_LDS_BYTES - Max LDS bytes per workgroup for TDM kernels (defaults to device max; K/M/G suffixes accepted)\n"); + printf(" TDM_TEMPORAL - TDM non-temporal cache hint (0=none, 1=loads, 2=stores, 3=both)\n"); + printf(" TDM_WORD_SIZE - TDM tensor descriptor data-size encoding (advanced; default 2)\n"); + printf(" NOTE: T/A (TDM / async load-store) executors do NOT honor GFX_* or\n"); + printf(" XCC_PREF_TABLE. Their threadblock size comes from TDM_BLOCK_SIZE.\n"); + printf(" USE_HIP_EVENTS - Use HIP events for GFX/DMA/TDM executor timing (0=CPU wall-clock)\n"); printf(" USE_HSA_DMA - Use hsa_amd_async_copy instead of hipMemcpy for non-targeted DMA execution\n"); printf(" USE_INTERACTIVE - Pause for user-input before starting transfer loop\n"); printf(" USE_SINGLE_STREAM - Use a single stream per GPU GFX executor instead of stream per Transfer\n"); @@ -443,9 +458,9 @@ class EnvVars { int numGpuDevices = TransferBench::GetNumExecutors(EXE_GPU_GFX); std::string nicSupport = ""; -#if NIC_EXEC_ENABLED - nicSupport = " (with NIC support)"; -#endif + if (IsIbvSymbolsReady()) { + nicSupport = " (with NIC support)"; + } if (!outputToCsv) { if (!hideEnv) printf("[Common] (Suppress by setting HIDE_ENV=1)\n"); } @@ -499,32 +514,31 @@ class EnvVars Print("GPU_MAX_HW_QUEUES", gpuMaxHwQueues, "Max %d hardware queues per GPU device", gpuMaxHwQueues); -#if NIC_EXEC_ENABLED - Print("IP_ADDRESS_FAMILY", ipAddressFamily, - "IP address family is set to IPv%d", ipAddressFamily); - - Print("IB_GID_INDEX", ibGidIndex, - "RoCE GID index is set to %s", (ibGidIndex < 0 ? "auto" : std::to_string(ibGidIndex).c_str())); - Print("IB_PORT_NUMBER", ibPort, - "IB port number is set to %d", ibPort); -#endif + if (IsIbvSymbolsReady()) { + Print("IP_ADDRESS_FAMILY", ipAddressFamily, + "IP address family is set to IPv%d", ipAddressFamily); + Print("IB_GID_INDEX", ibGidIndex, + "RoCE GID index is set to %s", (ibGidIndex < 0 ? "auto" : std::to_string(ibGidIndex).c_str())); + Print("IB_PORT_NUMBER", ibPort, + "IB port number is set to %d", ibPort); + Print("NIC_CHUNK_BYTES", nicChunkBytes, + "Sending %lu bytes at a time for NIC RDMA", nicChunkBytes); + Print("NIC_CQ_POLL_BATCH", nicCqPollBatch, + "Polling %d CQ entries per ibv_poll_cq call", nicCqPollBatch); + Print("NIC_RELAX_ORDER", nicRelaxedOrder, + "Using %s ordering for NIC RDMA", nicRelaxedOrder ? "relaxed" : "strict"); + Print("NIC_SERVICE_LEVEL", nicServiceLevel, + "IB service level (sl) set to %d", nicServiceLevel); + Print("NIC_TRAFFIC_CLASS", nicTrafficClass, + "RoCE traffic class (DSCP) set to %d", nicTrafficClass); + Print("ROCE_VERSION", roceVersion, + "RoCE version is set to %d", roceVersion); + } Print("MIN_VAR_SUBEXEC", minNumVarSubExec, "Using at least %d subexecutor(s) for variable subExec tranfers", minNumVarSubExec); Print("MAX_VAR_SUBEXEC", maxNumVarSubExec, "Using up to %s subexecutors for variable subExec transfers", maxNumVarSubExec ? std::to_string(maxNumVarSubExec).c_str() : "all available"); -#if NIC_EXEC_ENABLED - Print("NIC_CHUNK_BYTES", nicChunkBytes, - "Sending %lu bytes at a time for NIC RDMA", nicChunkBytes); - Print("NIC_CQ_POLL_BATCH", nicCqPollBatch, - "Polling %d CQ entries per ibv_poll_cq call", nicCqPollBatch); - Print("NIC_RELAX_ORDER", nicRelaxedOrder, - "Using %s ordering for NIC RDMA", nicRelaxedOrder ? "relaxed" : "strict"); - Print("NIC_SERVICE_LEVEL", nicServiceLevel, - "IB service level (sl) set to %d", nicServiceLevel); - Print("NIC_TRAFFIC_CLASS", nicTrafficClass, - "RoCE traffic class (DSCP) set to %d", nicTrafficClass); -#endif Print("NUM_ITERATIONS", numIterations, (numIterations == 0) ? "Running infinitely" : "Running %d %s", abs(numIterations), (numIterations > 0 ? " timed iteration(s)" : "seconds(s) per Test")); @@ -532,17 +546,28 @@ class EnvVars "Running %s subiterations", (numSubIterations == 0 ? "infinite" : std::to_string(numSubIterations)).c_str()); Print("NUM_WARMUPS", numWarmups, "Running %d warmup iteration(s) per Test", numWarmups); -#if NIC_EXEC_ENABLED - Print("ROCE_VERSION", roceVersion, - "RoCE version is set to %d", roceVersion); -#endif Print("SHOW_BORDERS", showBorders, "%s ASCII box-drawing characaters in tables", showBorders ? "Showing" : "Hiding"); Print("SHOW_ITERATIONS", showIterations, "%s per-iteration timing", showIterations ? "Showing" : "Hiding"); Print("SHOW_PERCENTILES", showPercentiles.empty() ? 0 : 1, "%s", showPercentiles.empty() ? "Disabled" : GetStr(showPercentiles).c_str()); + Print("TDM_BLOCK_ORDER", tdmBlockOrder, + "TDM worker ranking: %s", tdmBlockOrder == 0 ? "Sequential" : "Interleaved"); + Print("TDM_BLOCK_SIZE", tdmBlockSize, + "TDM threadblock size of %d", tdmBlockSize); + Print("TDM_MAX_LDS_BYTES", tdmMaxLdsBytes, + "%s", tdmMaxLdsBytes == INT_MAX + ? "Using device max LDS bytes per workgroup" + : (std::string("Capping LDS to ") + std::to_string(tdmMaxLdsBytes) + " bytes per workgroup").c_str()); + Print("TDM_TEMPORAL", tdmTemporal, + "%s", (tdmTemporal == 0 ? "Not using non-temporal hints" : + tdmTemporal == 1 ? "Using non-temporal load hint" : + tdmTemporal == 2 ? "Using non-temporal store hint" : + "Using non-temporal load and store hints")); + Print("TDM_WORD_SIZE", tdmWordSize, + "Using TDM tensor data-size encoding of %d", tdmWordSize); Print("USE_HIP_EVENTS", useHipEvents, - "Using %s for GFX/DMA Executor timing", useHipEvents ? "HIP events" : "CPU wall time"); + "Using %s for GFX/DMA/TDM Executor timing", useHipEvents ? "HIP events" : "CPU wall time"); Print("USE_HSA_DMA", useHsaDma, "Using %s for DMA execution", useHsaDma ? "hsa_amd_async_copy" : "hipMemcpyAsync"); Print("USE_INTERACTIVE", useInteractive, @@ -749,6 +774,13 @@ class EnvVars cfg.nic.trafficClass = nicTrafficClass; cfg.nic.roceVersion = roceVersion; + cfg.tdm.blockSize = tdmBlockSize; + cfg.tdm.maxLDSBytes = tdmMaxLdsBytes; + cfg.tdm.useHipEvents = useHipEvents; + cfg.tdm.blockOrder = tdmBlockOrder; + cfg.tdm.temporalMode = tdmTemporal; + cfg.tdm.wordSize = tdmWordSize; + return cfg; } }; diff --git a/src/client/Presets/Help.hpp b/src/client/Presets/Help.hpp index 26ede846..e80987d9 100644 --- a/src/client/Presets/Help.hpp +++ b/src/client/Presets/Help.hpp @@ -37,13 +37,16 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("# SRC 1 -> Executor -> DST 1\n"); printf("# SRC X DST Y\n"); printf("\n"); - printf("# Five Executors are supported by TransferBench\n"); - printf("# Executor: SubExecutor:\n"); - printf("# 1) CPU CPU thread\n"); - printf("# 2) GPU GPU threadblock/Compute Unit (CU)\n"); - printf("# 3) DMA N/A. (Must have single SRC, at least one DST)\n"); - printf("# 4) NIC Queue Pair\n"); - printf("# 5) Batched-DMA Batch item (Must have single SRC, at least one DST)\n"); + printf("# Eight Executors are supported by TransferBench\n"); + printf("# Executor: SubExecutor:\n"); + printf("# 1) CPU CPU thread\n"); + printf("# 2) GPU GPU threadblock/Compute Unit (CU)\n"); + printf("# 3) DMA N/A. (Must have single SRC, at least one DST)\n"); + printf("# 4) NIC Queue Pair\n"); + printf("# 5) Nearest NIC Queue Pair (uses closest NIC)\n"); + printf("# 6) Batched-DMA Batch item (Must have single SRC, at least one DST)\n"); + printf("# 7) TDM GPU threadblock/Compute Unit (CU)\n"); + printf("# 8) Async Load/Store GPU threadblock/Compute Unit (CU)\n"); printf("\n"); printf("# Each single line in the configuration file defines a set of Transfers (a Test) to run in parallel\n"); printf("\n"); @@ -71,6 +74,8 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("# - B: Batched-DMA-executor (Indexed from 0 to # GPUs - 1)\n"); printf("# - I#.#: NIC executor (Indexed from 0 to # NICs - 1)\n"); printf("# - N#.#: Nearest NIC executor (Indexed from 0 to # GPUs - 1)\n"); + printf("# - T: GPU TDM kernel (Indexed from 0 to # GPUs - 1)\n"); + printf("# - A: GPU async load/store (Indexed from 0 to # GPUs - 1)\n"); printf("# dstMemL : Destination memory locations (Where the data is to be written to)\n"); printf("# bytesL : Number of bytes to copy (0 means use command-line specified size)\n"); printf("# Must be a multiple of 4 and may be suffixed with ('K','M', or 'G')\n"); @@ -108,6 +113,12 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("## Single DMA executed Transfer between GPUs 0 and 1\n"); printf("1 1 (G0->D0->G1)\n"); printf("\n"); + printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for TDM path\n"); + printf("1 1 (G0->T0->G1)\n"); + printf("\n"); + printf("## Single GPU-executed Transfer between GPUs 0 and 1 using 1 CU for async load/store path\n"); + printf("1 1 (G0->A0->G1)\n"); + printf("\n"); printf("## Copy 1Mb from GPU0 to GPU1 with 4 CUs, and 2Mb from GPU1 to GPU0 with 8 CUs\n"); printf("-2 (G0->G0->G1 4 1M) (G1->G1->G0 8 2M)\n"); printf("\n"); diff --git a/src/client/Topology.hpp b/src/client/Topology.hpp index 09269377..1f5501eb 100644 --- a/src/client/Topology.hpp +++ b/src/client/Topology.hpp @@ -45,7 +45,7 @@ static int RemappedCpuIndex(int origIdx) static void PrintNicToGPUTopo(bool outputToCsv) { -#ifdef NIC_EXEC_ENABLED + if (!IsIbvSymbolsReady()) return; printf(" NIC | Device Name | Active | PCIe Bus ID | NUMA | Closest GPU(s) | GID Index | GID Descriptor\n"); if(!outputToCsv) printf("-----+-------------+--------+--------------+------+----------------+-----------+-------------------\n"); @@ -73,7 +73,6 @@ static void PrintNicToGPUTopo(bool outputToCsv) ); } printf("\n"); -#endif } void DisplaySingleRankTopology(bool outputToCsv) diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index 9730ace6..af9a1696 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -439,6 +439,8 @@ namespace TransferBench::Utils case EXE_NIC: return "NIC"; case EXE_NIC_NEAREST: return "NIC"; case EXE_GPU_BDMA: return "BMA"; + case EXE_GPU_TDM: return "TDM"; + case EXE_GPU_ALS: return "ALS"; default: return "N/A"; } } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 5f22f52c..b69c1812 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -25,6 +25,9 @@ THE SOFTWARE. #include #include #include +#include +#include +#include #include #include #include @@ -55,9 +58,7 @@ THE SOFTWARE. #include #include -#ifdef NIC_EXEC_ENABLED -#include -#endif +#include "IbvDynLoad.hpp" #ifdef MPI_COMM_ENABLED #include @@ -77,9 +78,9 @@ THE SOFTWARE. #ifdef AMD_SMI_ENABLED #include "amd_smi/amdsmi.h" #endif +#include "tdmCopy.h" #endif /// @endcond - // Batched DMA executor is only supported with HIP >= 7.1 and CUDA 12.8 #if (defined(HIP_VERSION) && (HIP_VERSION >= 70100000)) || (defined(CUDA_VERSION) && (CUDA_VERSION >= 12080)) #define BMA_EXEC_ENABLED @@ -107,10 +108,15 @@ namespace TransferBench EXE_NIC = 3, ///< NIC RDMA executor (subExecutor = queue pair) EXE_NIC_NEAREST = 4, ///< NIC RDMA nearest executor (subExecutor = queue pair) EXE_GPU_BDMA = 5, ///< GPU Batched SDMA executor (subExecutor = batch item) + EXE_GPU_TDM = 6, ///< GPU TDM executor (subExecutor = threadblock/CU) + EXE_GPU_ALS = 7, ///< GPU async load/store (subExecutor — reserved; stub uses 1-D launch) }; - char const ExeTypeStr[7] = "CGDINB"; + char const ExeTypeStr[9] = "CGDINBTA"; inline bool IsCpuExeType(ExeType e){ return e == EXE_CPU; } - inline bool IsGpuExeType(ExeType e){ return e == EXE_GPU_GFX || e == EXE_GPU_DMA || e == EXE_GPU_BDMA; } + inline bool IsGpuExeType(ExeType e){ + return e == EXE_GPU_GFX || e == EXE_GPU_DMA || e == EXE_GPU_BDMA + || e == EXE_GPU_TDM || e == EXE_GPU_ALS; + } inline bool IsNicExeType(ExeType e){ return e == EXE_NIC || e == EXE_NIC_NEAREST; } /** @@ -275,6 +281,17 @@ namespace TransferBench int useNuma = 0; ///< Switch to closest numa thread for execution }; + struct TdmOptions + { + int blockSize = 256; ///< Size of each threadblock (must be multiple of 32) + int maxLDSBytes = INT_MAX; ///< Maximum number of bytes of __shared__ memory to use + int useHipEvents = 1; ///< Use HIP events for timing the TDM executor (0=CPU wall-clock) + + int wordSize = 2; ///< (unused for now: TdmCopy kernel to expand) Tensor descriptor data-size encoding (group1.dataSize); advanced + int temporalMode = 0; ///< (unused for now: TdmCopy kernel to expand) Non-temporal cache hint (0=none, 1=loads, 2=stores, 3=both) + int blockOrder = 0; ///< (unused for now: only one transfer per stream) Worker ranking across blocks + }; + /** * Configuration options for performing Transfers @@ -287,6 +304,7 @@ namespace TransferBench GfxOptions gfx; ///< GFX executor options DmaOptions dma; ///< DMA executor options NicOptions nic; ///< NIC executor options + TdmOptions tdm; ///< TDM executor options }; /** @@ -620,6 +638,7 @@ namespace TransferBench // Enumerations #define hipDeviceAttributeClockRate cudaDevAttrClockRate #define hipDeviceAttributeMultiprocessorCount cudaDevAttrMultiProcessorCount + #define hipDeviceAttributeMaxSharedMemoryPerBlock cudaDevAttrMaxSharedMemoryPerBlock #define hipDeviceAttributeWarpSize cudaDevAttrWarpSize #define hipErrorPeerAccessAlreadyEnabled cudaErrorPeerAccessAlreadyEnabled #define hipFuncCachePreferShared cudaFuncCachePreferShared @@ -649,6 +668,7 @@ namespace TransferBench #define hipGetDeviceCount cudaGetDeviceCount #define hipGetDeviceProperties cudaGetDeviceProperties #define hipGetErrorString cudaGetErrorString + #define hipGetLastError cudaGetLastError #define hipHostFree cudaFreeHost #define hipHostMalloc cudaMallocHost #define hipMalloc cudaMalloc @@ -1669,16 +1689,34 @@ namespace { return ERR_NONE; } -#if defined(NIC_EXEC_ENABLED) && defined(HAVE_DMABUF_SUPPORT) && !defined(__NVCC__) +#if defined(__NVCC__) + static bool CheckDmabufSupport() + { + return false; + } + static ErrResult ExportDmabuf(void* gpuPtr, size_t numBytes, int& dmabufFd, uint64_t& dmabufOffset) + { + return {ERR_FATAL, "DMA-BUF export not yet supported on NVIDIA platform"}; + } +#else + hsa_status_t (*pfn_hsa_amd_portable_export_dmabuf)(const void*, size_t, int*, uint64_t*); // Check kernel configuration for required DMA-BUF support // Returns true if kernel supports CONFIG_DMABUF_MOVE_NOTIFY and CONFIG_PCI_P2PDMA - static bool CheckKernelDmabufSupport() + static bool CheckDmabufSupport() { static int support = -1; // -1: not checked, 0: disabled, 1: enabled if (support != -1) { return support; } + // Check hsa_amd_portable_export_dmabuf and ibv_reg_dmabuf_mr symbols are available + // rocr and hsa_ext_amd header is always mandatory, so no need to check for them + pfn_hsa_amd_portable_export_dmabuf = + (hsa_status_t (*)(const void*, size_t, int*, uint64_t*))dlsym(RTLD_DEFAULT, "hsa_amd_portable_export_dmabuf"); + if (pfn_hsa_amd_portable_export_dmabuf == nullptr || !IsIbvDmabufPresent()) { + support = 0; + return support; + } struct utsname utsname; FILE* fp = NULL; @@ -1802,7 +1840,7 @@ namespace { // Export the aligned GPU buffer as DMA-BUF uint64_t exportOffset = 0; - hsa_status_t status = hsa_amd_portable_export_dmabuf(alignedPtr, alignedSize, &dmabufFd, &exportOffset); + hsa_status_t status = pfn_hsa_amd_portable_export_dmabuf(alignedPtr, alignedSize, &dmabufFd, &exportOffset); if (status != HSA_STATUS_SUCCESS) { return {ERR_FATAL, "Failed to export DMA-BUF: hsa_amd_portable_export_dmabuf returned %d", status}; @@ -2109,16 +2147,51 @@ namespace { } } - // Check NIC options -#ifdef NIC_EXEC_ENABLED - if (cfg.nic.chunkBytes == 0 || (cfg.nic.chunkBytes % 4 != 0)) { - errors.push_back({ERR_FATAL, "[nic.chunkBytes] must be a non-negative multiple of 4"}); - } - if (cfg.nic.cqPollBatch <= 0) { - errors.push_back({ERR_FATAL, "[nic.cqPollBatch] must be positive"}); + // Check TDM options (TDM is AMD-only; on NVIDIA the executor is a no-op stub) +#if !defined(__NVCC__) + if (cfg.tdm.blockSize <= 0 || cfg.tdm.blockSize % 32 || cfg.tdm.blockSize > MAX_BLOCKSIZE) + errors.push_back({ERR_FATAL, + "[tdm.blockSize] must be a positive multiple of 32 less than or equal to %d", + MAX_BLOCKSIZE}); + + if (cfg.tdm.temporalMode < 0 || cfg.tdm.temporalMode > 3) + errors.push_back({ERR_FATAL, + "[tdm.temporalMode] must be non-negative and less than or equal to 3"}); +/* + // TDM only exercises the descriptor data-size encoding of 2 (matches its float buffers) for now. + if (cfg.tdm.wordSize != 2) + errors.push_back({ERR_FATAL, + "[tdm.wordSize] must be 2 (the only descriptor data-size encoding TDM currently uses)"}); +*/ + if (cfg.tdm.maxLDSBytes <= 0) + errors.push_back({ERR_FATAL, "[tdm.maxLDSBytes] must be positive"}); + else if (cfg.tdm.maxLDSBytes != INT_MAX) { + // Runtime clamps to the device max (see RunTdmExecutor); warn if the user's cap exceeds it. + int const numGpus = GetNumExecutors(EXE_GPU_TDM); + int minDeviceMax = INT_MAX; + for (int i = 0; i < numGpus; i++) { + int deviceMax = 0; + if (hipDeviceGetAttribute(&deviceMax, hipDeviceAttributeMaxSharedMemoryPerBlock, i) == hipSuccess + && deviceMax > 0) + minDeviceMax = std::min(minDeviceMax, deviceMax); + } + if (minDeviceMax != INT_MAX && cfg.tdm.maxLDSBytes > minDeviceMax) + errors.push_back({ERR_WARN, + "[tdm.maxLDSBytes] (%d) exceeds device max shared memory per block (%d); will be clamped", + cfg.tdm.maxLDSBytes, minDeviceMax}); } #endif + // Check NIC options + if (IsIbvSymbolsReady()) { + if (cfg.nic.chunkBytes == 0 || (cfg.nic.chunkBytes % 4 != 0)) { + errors.push_back({ERR_FATAL, "[nic.chunkBytes] must be a non-negative multiple of 4"}); + } + if (cfg.nic.cqPollBatch <= 0) { + errors.push_back({ERR_FATAL, "[nic.cqPollBatch] must be positive"}); + } + } + // NVIDIA specific #if defined(__NVCC__) if (cfg.data.validateDirect) @@ -2242,7 +2315,7 @@ namespace { // Each subexecutor is assigned a multiple of cfg.data.blockBytes, however this may // mean that some subexecutors might not have any work assigned to them if the amount to // transfer is small - if (t.exeDevice.exeType == EXE_GPU_GFX || t.exeDevice.exeType == EXE_CPU || t.exeDevice.exeType == EXE_GPU_BDMA) { + if (IsGpuExeType(t.exeDevice.exeType)) { size_t const N = t.numBytes / sizeof(float); int const targetMultiple = cfg.data.blockBytes / sizeof(float); int const maxSubExecToUse = std::min((size_t)(N + targetMultiple - 1) / targetMultiple, @@ -2300,6 +2373,44 @@ namespace { hasFatalError = true; } break; + case EXE_GPU_ALS: + errors.push_back({ERR_FATAL, + "Transfer %d: Async load/store executor (A) is not supported yet", i}); + hasFatalError = true; + break; + case EXE_GPU_TDM: + if (t.srcs.size() != 1 || t.dsts.size() != 1) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel currently requires exactly 1 SRC and 1 DST", i}); + hasFatalError = true; + break; + } + if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel: device index must be between 0 and %d (instead of %d) for rank %d", + i, numExecutors - 1, t.exeDevice.exeIndex, t.exeDevice.exeRank}); + hasFatalError = true; + break; + } + if (t.exeSubIndex != -1) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM executor does not support subindices yet", i}); + hasFatalError = true; + break; + } +#if defined(__NVCC__) + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel is not supported on NVIDIA hardware", i}); + hasFatalError = true; +#else + if (!tdm::IsTdmCopySupported(t.exeDevice.exeIndex)) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel requires gfx1250 hardware, but GPU %d is not supported", + i, t.exeDevice.exeIndex}); + hasFatalError = true; + } +#endif + break; case EXE_GPU_GFX: if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { errors.push_back({ERR_FATAL, @@ -2320,7 +2431,6 @@ namespace { errors.push_back({ERR_FATAL, "Transfer %d: GFX subIndex (XCC) must be between 0 and %d for rank %d", i, numSubIndices - 1, t.exeDevice.exeRank}); hasFatalError = true; - break; } #endif } @@ -2488,7 +2598,7 @@ namespace { break; #endif case EXE_NIC: case EXE_NIC_NEAREST: -#ifdef NIC_EXEC_ENABLED + if (IsIbvSymbolsReady()) { // NIC Executors can only execute a copy operation if (t.srcs.size() != 1 || t.dsts.size() != 1) { @@ -2542,11 +2652,10 @@ namespace { hasFatalError = true; break; } + } else { + errors.push_back({ERR_FATAL, "Transfer %d: NIC executor is requested but is not available.", i}); + hasFatalError = true; } -#else - errors.push_back({ERR_FATAL, "Transfer %d: NIC executor is requested but is not available.", i}); - hasFatalError = true; -#endif break; } @@ -2671,6 +2780,22 @@ namespace { } break; } + case EXE_GPU_TDM: +// case EXE_GPU_ALS: + { + int numGpuSubExec = GetNumSubExecutors(exeDevice); + if (totalSubExecs[exeDevice] > numGpuSubExec) + errors.push_back({ERR_WARN, + "GPU %d TDM requests %d total subexecutors however only %d CUs available. " + "Serialization may occur", + exeDevice.exeIndex, totalSubExecs[exeDevice], numGpuSubExec}); + if (transferCount[exeDevice] > gpuMaxHwQueues) { + errors.push_back({ERR_WARN, + "GPU %d TDM attempting %d parallel transfers, however GPU_MAX_HW_QUEUES only set to %d", + exeDevice.exeIndex, transferCount[exeDevice], gpuMaxHwQueues}); + } + break; + } case EXE_GPU_DMA: { // Check that if executor subindices are used, all Transfers specify executor subindices @@ -2767,7 +2892,6 @@ namespace { #endif // For IBV executor -#ifdef NIC_EXEC_ENABLED int srcNicIndex; ///< SRC NIC index int dstNicIndex; ///< DST NIC index ibv_context* srcContext; ///< Device context for SRC NIC @@ -2792,7 +2916,6 @@ namespace { bool srcIsExeNic; ///< Whether SRC or DST NIC initiates traffic vector> sgePerQueuePair; ///< Scatter-gather elements per queue pair vector>sendWorkRequests; ///< Send work requests per queue pair -#endif // For BMA executor #ifdef BMA_EXEC_ENABLED @@ -2850,7 +2973,6 @@ namespace { } }; -#ifdef NIC_EXEC_ENABLED // Structure to track information about IBV devices struct IbvDevice { @@ -2863,12 +2985,11 @@ namespace { std::string gidDescriptor; bool isRoce; }; -#endif -#ifdef NIC_EXEC_ENABLED // Function to collect information about IBV devices //======================================================================================== -static bool IsConfiguredGid(union ibv_gid const& gid) + + static bool IsConfiguredGid(union ibv_gid const& gid) { const struct in6_addr *a = (struct in6_addr *) gid.raw; int trailer = (a->s6_addr32[1] | a->s6_addr32[2] | a->s6_addr32[3]); @@ -2981,7 +3102,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) static vector ibvDeviceList = {}; // Build list on first use - if (!isInitialized) { + if (IsIbvSymbolsReady() && !isInitialized) { // Query the number of IBV devices int numIbvDevices = 0; @@ -3066,9 +3187,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } return ibvDeviceList; } -#endif // NIC_EXEC_ENABLED -#ifdef NIC_EXEC_ENABLED // PCIe-related functions //======================================================================================== @@ -3253,9 +3372,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } return matches; } -#endif // NIC_EXEC_ENABLED -#ifdef NIC_EXEC_ENABLED // IB Verbs-related functions //======================================================================================== @@ -3426,16 +3543,12 @@ static bool IsConfiguredGid(union ibv_gid const& gid) if (!dmabufStatusPrinted) { dmabufStatusPrinted = true; printf("[INFO] Rank %d DMA-BUF support: ", GetRank()); -#if defined(HAVE_DMABUF_SUPPORT) && !defined(__NVCC__) - bool kernelSupport = CheckKernelDmabufSupport(); + bool kernelSupport = CheckDmabufSupport(); if (kernelSupport) { printf("ENABLED\n"); } else { - printf("DISABLED (kernel config missing, using standard ibv_reg_mr)\n"); + printf("DISABLED (kernel config or export symbol missing, using standard ibv_reg_mr)\n"); } -#else - printf("DISABLED (using standard ibv_reg_mr)\n"); -#endif } } @@ -3464,27 +3577,22 @@ static bool IsConfiguredGid(union ibv_gid const& gid) IBV_PTR_CALL(rss.srcProtect, ibv_alloc_pd, rss.srcContext); // Export DMA-BUF for SRC memory if it's GPU memory -#if defined(HAVE_DMABUF_SUPPORT) && !defined(__NVCC__) - if (!t.srcs.empty() && IsGpuMemType(t.srcs[0].memType) && CheckKernelDmabufSupport()) { + if (CheckDmabufSupport() && !t.srcs.empty() && IsGpuMemType(t.srcs[0].memType)) { ERR_CHECK(ExportDmabuf(rss.srcMem[0], rss.numBytes, rss.srcDmabufFd, rss.srcDmabufOffset)); if (System::Get().IsVerbose()) { printf("[INFO] Rank %d exported SRC GPU memory as DMA-BUF (fd=%d, offset=%lu)\n", GetRank(), rss.srcDmabufFd, rss.srcDmabufOffset); } } -#endif // Register SRC memory region -#ifdef HAVE_DMABUF_SUPPORT if (rss.srcDmabufFd >= 0) { IBV_PTR_CALL(rss.srcMemRegion, ibv_reg_dmabuf_mr, rss.srcProtect, rss.srcDmabufOffset, rss.numBytes, (uint64_t)rss.srcMem[0], rss.srcDmabufFd, rdmaMemRegFlags); if (System::Get().IsVerbose()) { printf("[INFO] Rank %d registered SRC memory using ibv_reg_dmabuf_mr\n", GetRank()); } - } else -#endif - { + } else { IBV_PTR_CALL(rss.srcMemRegion, ibv_reg_mr, rss.srcProtect, rss.srcMem[0], rss.numBytes, rdmaMemRegFlags); if (System::Get().IsVerbose()) { printf("[INFO] Rank %d registered SRC memory using ibv_reg_mr (standard path)\n", GetRank()); @@ -3530,27 +3638,22 @@ static bool IsConfiguredGid(union ibv_gid const& gid) IBV_PTR_CALL(rss.dstProtect, ibv_alloc_pd, rss.dstContext); // Export DMA-BUF for DST memory if it's GPU memory -#if defined(HAVE_DMABUF_SUPPORT) && !defined(__NVCC__) - if (!t.dsts.empty() && IsGpuMemType(t.dsts[0].memType) && CheckKernelDmabufSupport()) { + if (CheckDmabufSupport() && !t.dsts.empty() && IsGpuMemType(t.dsts[0].memType)) { ERR_CHECK(ExportDmabuf(rss.dstMem[0], rss.numBytes, rss.dstDmabufFd, rss.dstDmabufOffset)); if (System::Get().IsVerbose()) { printf("[INFO] Rank %d exported DST GPU memory as DMA-BUF (fd=%d, offset=%lu)\n", GetRank(), rss.dstDmabufFd, rss.dstDmabufOffset); } } -#endif // Register DST memory region -#ifdef HAVE_DMABUF_SUPPORT if (rss.dstDmabufFd >= 0) { IBV_PTR_CALL(rss.dstMemRegion, ibv_reg_dmabuf_mr, rss.dstProtect, rss.dstDmabufOffset, rss.numBytes, (uint64_t)rss.dstMem[0], rss.dstDmabufFd, rdmaMemRegFlags); if (System::Get().IsVerbose()) { printf("[INFO] Rank %d registered DST memory using ibv_reg_dmabuf_mr\n", GetRank()); } - } else -#endif - { + } else { IBV_PTR_CALL(rss.dstMemRegion, ibv_reg_mr, rss.dstProtect, rss.dstMem[0], rss.numBytes, rdmaMemRegFlags); if (System::Get().IsVerbose()) { printf("[INFO] Rank %d registered DST memory using ibv_reg_mr (standard path)\n", GetRank()); @@ -3721,16 +3824,16 @@ static bool IsConfiguredGid(union ibv_gid const& gid) if (isDstRank) IBV_CALL(ibv_dereg_mr, rss.dstMemRegion); // Close DMA-BUF file descriptors -#if defined(HAVE_DMABUF_SUPPORT) && !defined(__NVCC__) - if (isSrcRank && rss.srcDmabufFd >= 0) { - close(rss.srcDmabufFd); - rss.srcDmabufFd = -1; - } - if (isDstRank && rss.dstDmabufFd >= 0) { - close(rss.dstDmabufFd); - rss.dstDmabufFd = -1; + if (CheckDmabufSupport()) { + if (isSrcRank && rss.srcDmabufFd >= 0) { + close(rss.srcDmabufFd); + rss.srcDmabufFd = -1; + } + if (isDstRank && rss.dstDmabufFd >= 0) { + close(rss.dstDmabufFd); + rss.dstDmabufFd = -1; + } } -#endif // Destroy queue pairs if (isSrcRank) { @@ -3758,7 +3861,6 @@ static bool IsConfiguredGid(union ibv_gid const& gid) return ERR_NONE; } -#endif // NIC_EXEC_ENABLED // Data validation-related functions //======================================================================================== @@ -4327,13 +4429,14 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } // Prepare additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) - && exeDevice.exeRank == localRank) { + if (IsGpuExeType(exeDevice.exeType) && exeDevice.exeRank == localRank) { ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); // Determine how many streams to use int const numStreamsToUse = (exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || - (exeDevice.exeType == EXE_GPU_GFX && cfg.gfx.useMultiStream)) + (exeDevice.exeType == EXE_GPU_GFX && cfg.gfx.useMultiStream) || + exeDevice.exeType == EXE_GPU_TDM || + exeDevice.exeType == EXE_GPU_ALS) ? exeInfo.resources.size() : 1; exeInfo.streams.resize(numStreamsToUse); @@ -4351,7 +4454,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { + if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents || cfg.tdm.useHipEvents) { exeInfo.startEvents.resize(numStreamsToUse); exeInfo.stopEvents.resize(numStreamsToUse); for (int i = 0; i < numStreamsToUse; ++i) { @@ -4361,8 +4464,9 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - // Prepare for GPU GFX executor - if (exeDevice.exeType == EXE_GPU_GFX && exeDevice.exeRank == localRank) { + // Prepare for GPU GFX / TDM executor (both consume SubExecParam from GPU memory) + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_TDM) && + exeDevice.exeRank == localRank) { // Allocate one contiguous chunk of GPU memory for threadblock parameters // This allows support for executing one transfer per stream, or all transfers in a single stream #if !defined(__NVCC__) @@ -4383,7 +4487,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) exeDevice.exeIndex)); #endif int transferOffset = 0; - if (cfg.gfx.useMultiStream || cfg.gfx.blockOrder == 0) { + // TDM always runs one transfer per stream, for now) + if (exeDevice.exeType == EXE_GPU_TDM || cfg.gfx.useMultiStream || cfg.gfx.blockOrder == 0) { // Threadblocks are ordered sequentially one transfer at a time for (auto& rss : exeInfo.resources) { rss.subExecParamGpuPtr = exeInfo.subExecParamGpu + transferOffset; @@ -4445,18 +4550,19 @@ static bool IsConfiguredGid(union ibv_gid const& gid) // Prepare for NIC-based executors if (IsNicExeType(exeDevice.exeType)) { -#ifdef NIC_EXEC_ENABLED - for (auto& rss : exeInfo.resources) { - Transfer const& t = transfers[rss.transferIdx]; - ERR_CHECK(PrepareNicTransferResources(cfg, exeDevice, t, rss)); + if (IsIbvSymbolsReady()) { + for (auto& rss : exeInfo.resources) { + Transfer const& t = transfers[rss.transferIdx]; + ERR_CHECK(PrepareNicTransferResources(cfg, exeDevice, t, rss)); + } + } else { + return {ERR_FATAL, "RDMA executor is not supported"}; } -#else - return {ERR_FATAL, "RDMA executor is not supported"}; -#endif } - // Check that GPU wallclock rate is non-zero - if (exeDevice.exeType == EXE_GPU_GFX && exeInfo.wallClockRate == 0 && exeDevice.exeRank == localRank) { + // Check that GPU wallclock rate is non-zero (GFX and TDM both use it for in-kernel cycle timing) + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_TDM) && + exeInfo.wallClockRate == 0 && exeDevice.exeRank == localRank) { if (getenv("TB_WALLCLOCK_RATE")) { exeInfo.wallClockRate = atoi(getenv("TB_WALLCLOCK_RATE")); return {ERR_WARN, @@ -4549,27 +4655,23 @@ static bool IsConfiguredGid(union ibv_gid const& gid) #endif // Destroy NIC related resources -#ifdef NIC_EXEC_ENABLED - if (IsNicExeType(exeDevice.exeType)) { + if (IsIbvSymbolsReady() && IsNicExeType(exeDevice.exeType)) { ERR_CHECK(TeardownNicTransferResources(rss, t)); } -#endif } // Teardown additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) - && exeDevice.exeRank == localRank) { + if (IsGpuExeType(exeDevice.exeType) && exeDevice.exeRank == localRank) { for (auto stream : exeInfo.streams) ERR_CHECK(hipStreamDestroy(stream)); - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { - for (auto event : exeInfo.startEvents) - ERR_CHECK(hipEventDestroy(event)); - for (auto event : exeInfo.stopEvents) - ERR_CHECK(hipEventDestroy(event)); - } + for (auto event : exeInfo.startEvents) + ERR_CHECK(hipEventDestroy(event)); + for (auto event : exeInfo.stopEvents) + ERR_CHECK(hipEventDestroy(event)); } - if (exeDevice.exeType == EXE_GPU_GFX && exeDevice.exeRank == localRank) { + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_TDM) && + exeDevice.exeRank == localRank) { #if !defined(__NVCC__) MemType memType = MEM_GPU; #else @@ -4679,7 +4781,6 @@ static bool IsConfiguredGid(union ibv_gid const& gid) return ERR_NONE; } -#ifdef NIC_EXEC_ENABLED // Execution of a single NIC Transfer static ErrResult ExecuteNicTransfer(int const iteration, ConfigOptions const& cfg, @@ -4780,7 +4881,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } return ERR_NONE; } -#endif + // GFX Executor-related functions //======================================================================================== @@ -5703,6 +5804,161 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } #endif // BMA_EXEC_ENABLED +// TDM Executor-related functions (gfx1250+ only) +//======================================================================================== +#if TDM_SUPPORTED + __global__ void GpuTdmKernel(SubExecParam* params, + size_t ldsBytes, + int dataSize, + int temporalMode, + int numSubIterations) + { + // Retained for now but unused: tdmCopy fixes the descriptor data-size (4-byte word) + // and cache policy internally. + (void)dataSize; + (void)temporalMode; + + int64_t startCycle; + bool const shouldRecordTiming = (threadIdx.x == 0); + if (shouldRecordTiming) startCycle = GetTimestamp(); + + extern __shared__ __align__(128) float shmem[]; + + // Each threadblock is a subexecutor (mirrors GpuCopyKernel). + int const subExecIdx = (int)blockIdx.x; + + SubExecParam& p = params[subExecIdx]; + if (p.N == 0) return; + + float const* __restrict__ src = (float const*)p.src[0]; + float* __restrict__ dst = (float*)p.dst[0]; + + // Block-collective copy: hand the whole dynamic LDS allocation (ldsBytes) to tdmCopy, + // which partitions both the work and the staging buffer across every warp in the block + // and handles head/tail alignment internally. + size_t const sizeBytes = p.N * sizeof(float); + + int subIterations = 0; + while (1) { + tdm::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes); + __syncthreads(); + if (++subIterations == numSubIterations) break; + } + + if (shouldRecordTiming) { + p.stopCycle = GetTimestamp(); + p.startCycle = startCycle; + GetHwId(p.hwId); + GetXccId(p.xccId); + } + } +#else + // gfx1250 tensor TDM builtins unavailable for this translation: emit empty kernel stubs with + // the exact launch signatures so the host-side launch path still links. They are never + // dispatched on non-gfx1250 or nvidia hardware (see TransfersHaveErrors). + __global__ void GpuTdmKernel(SubExecParam*, size_t, int, int, int) {} +#endif // TDM_SUPPORTED + + static ErrResult ExecuteTdmTransfer(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + hipStream_t const stream, + hipEvent_t const startEvent, + hipEvent_t const stopEvent, + int const numSubExecs, + TransferResources& rss) + { + // Block-collective LDS budget handed straight to tdmCopy: the whole per-block dynamic + // shared-memory allocation, clamped to the device max (and cfg.tdm.maxLDSBytes). tdmCopy + // partitions this pool across the block's warps internally, so no per-tile sizing is needed. + int maxShmem = 0; + ERR_CHECK(hipDeviceGetAttribute(&maxShmem, hipDeviceAttributeMaxSharedMemoryPerBlock, exeIndex)); + maxShmem = std::min(maxShmem, cfg.tdm.maxLDSBytes); + size_t const ldsBytes = static_cast(maxShmem); + + dim3 const gridSize(numSubExecs); + dim3 const blockSize(cfg.tdm.blockSize); + + auto cpuStart = std::chrono::high_resolution_clock::now(); + + // Sub-iterations run on-device inside a single launch (matches the GFX GpuCopyKernel path), so + // kernel-launch and stream-sync overhead is not folded into the per-sub-iteration timing. + if (startEvent) + ERR_CHECK(hipEventRecord(startEvent, stream)); +#if defined(__NVCC__) + GpuTdmKernel<<>>(rss.subExecParamGpuPtr, + ldsBytes, cfg.tdm.wordSize, cfg.tdm.temporalMode, cfg.general.numSubIterations); +#else + hipLaunchKernelGGL(GpuTdmKernel, gridSize, blockSize, + ldsBytes, stream, rss.subExecParamGpuPtr, + ldsBytes, cfg.tdm.wordSize, cfg.tdm.temporalMode, cfg.general.numSubIterations); +#endif + ERR_CHECK(hipGetLastError()); + if (stopEvent) + ERR_CHECK(hipEventRecord(stopEvent, stream)); + ERR_CHECK(hipStreamSynchronize(stream)); + + // Per-transfer timing (mirrors ExecuteGpuTransfer's multistream branch): CPU wall-clock baseline, + // overridden by HIP events when enabled. CU set is read from the per-block SubExecParam slices. + if (iteration >= 0) { + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 + / cfg.general.numSubIterations; + if (startEvent && stopEvent) { + float gpuDeltaMsec; + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); + deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations; + } + rss.totalDurationMsec += deltaMsec; + if (cfg.general.recordPerIteration) { + rss.perIterMsec.push_back(deltaMsec); + std::set> CUs; + for (int i = 0; i < numSubExecs; i++) { + CUs.insert(std::make_pair(rss.subExecParamGpuPtr[i].xccId, + GetId(rss.subExecParamGpuPtr[i].hwId))); + } + rss.perIterCUs.push_back(CUs); + } + } + return ERR_NONE; + } + + static ErrResult RunTdmExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + ExeInfo& exeInfo) + { + auto cpuStart = std::chrono::high_resolution_clock::now(); + ERR_CHECK(hipSetDevice(exeIndex)); + + vector> asyncTransfers; + for (size_t i = 0; i < exeInfo.resources.size(); i++) { + hipEvent_t startEv = (cfg.tdm.useHipEvents && !exeInfo.startEvents.empty()) ? exeInfo.startEvents[i] : nullptr; + hipEvent_t stopEv = (cfg.tdm.useHipEvents && !exeInfo.stopEvents.empty()) ? exeInfo.stopEvents[i] : nullptr; + int const numSubExecs = + static_cast(exeInfo.resources[i].subExecParamCpu.size()); + asyncTransfers.emplace_back(std::async(std::launch::async, + ExecuteTdmTransfer, + iteration, + std::cref(cfg), + exeIndex, + exeInfo.streams[i], + startEv, + stopEv, + numSubExecs, + std::ref(exeInfo.resources[i]))); + } + for (auto& asyncTransfer : asyncTransfers) + ERR_CHECK(asyncTransfer.get()); + + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 + / cfg.general.numSubIterations; + if (iteration >= 0) + exeInfo.totalDurationMsec += deltaMsec; + return ERR_NONE; + } + // Executor-related functions //======================================================================================== static ErrResult RunExecutor(int const iteration, @@ -5713,10 +5969,9 @@ static bool IsConfiguredGid(union ibv_gid const& gid) switch (exeDevice.exeType) { case EXE_CPU: return RunCpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); case EXE_GPU_GFX: return RunGpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_TDM: return RunTdmExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); case EXE_GPU_DMA: return RunDmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); -#ifdef NIC_EXEC_ENABLED case EXE_NIC: return RunNicExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); -#endif #ifdef BMA_EXEC_ENABLED case EXE_GPU_BDMA: return RunBmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); #endif @@ -6123,11 +6378,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) TransferResult& tfrResult = results.tfrResults[transferIdx]; tfrResult.exeDevice = exeDevice; -#ifdef NIC_EXEC_ENABLED tfrResult.exeDstDevice = {exeDevice.exeType, rss.dstNicIndex}; -#else - tfrResult.exeDstDevice = exeDevice; -#endif tfrResult.numBytes = rss.numBytes; tfrResult.avgDurationMsec = rss.totalDurationMsec / numTimedIterations; tfrResult.avgBandwidthGbPerSec = (rss.numBytes / 1.0e6) / tfrResult.avgDurationMsec; @@ -6413,6 +6664,12 @@ static bool IsConfiguredGid(union ibv_gid const& gid) result |= RecursiveWildcardTransferExpansion(wc, baseRankIndex, numBytes, numSubExecs, transfers); wc.exe.exeSubIndices[0] = -2; return result; + case EXE_GPU_TDM: + case EXE_GPU_ALS: + wc.exe.exeSubIndices[0] = -1; + result |= RecursiveWildcardTransferExpansion(wc, baseRankIndex, numBytes, numSubExecs, transfers); + wc.exe.exeSubIndices[0] = -2; + return result; case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: { // Iterate over all available subindices @@ -7297,9 +7554,11 @@ static bool IsConfiguredGid(union ibv_gid const& gid) int numGpus = 0; hipError_t status = hipGetDeviceCount(&numGpus); if (status != hipSuccess) numGpus = 0; - topo.numExecutors[EXE_GPU_GFX] = numGpus; - topo.numExecutors[EXE_GPU_DMA] = numGpus; + topo.numExecutors[EXE_GPU_GFX] = numGpus; + topo.numExecutors[EXE_GPU_DMA] = numGpus; topo.numExecutors[EXE_GPU_BDMA] = numGpus; + topo.numExecutors[EXE_GPU_TDM] = numGpus; + topo.numExecutors[EXE_GPU_ALS] = numGpus; std::vector gpuArchNames(numGpus); @@ -7323,6 +7582,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) topo.executorName[{EXE_GPU_GFX, exeIndex}] = gpuName; topo.executorName[{EXE_GPU_DMA, exeIndex}] = gpuName; topo.executorName[{EXE_GPU_BDMA, exeIndex}] = gpuName; + topo.executorName[{EXE_GPU_TDM, exeIndex}] = gpuName + " [TDM]"; + topo.executorName[{EXE_GPU_ALS, exeIndex}] = gpuName + " [async LD/ST]"; #if !defined(__NVCC__) hsa_agent_t gpuAgent = gpuAgents[exeIndex]; @@ -7352,9 +7613,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid) topo.numExecutorSubIndices[{EXE_GPU_GFX, exeIndex}] = numXccs; topo.numExecutorSubIndices[{EXE_GPU_DMA, exeIndex}] = numDmaEngines; topo.numExecutorSubIndices[{EXE_GPU_BDMA, exeIndex}] = 0; + topo.numExecutorSubIndices[{EXE_GPU_TDM, exeIndex}] = 0; + topo.numExecutorSubIndices[{EXE_GPU_ALS, exeIndex}] = 0; topo.numSubExecutors[{EXE_GPU_GFX, exeIndex}] = numDeviceCUs; topo.numSubExecutors[{EXE_GPU_DMA, exeIndex}] = 1; topo.numSubExecutors[{EXE_GPU_BDMA, exeIndex}] = numDmaEngines; + topo.numSubExecutors[{EXE_GPU_TDM, exeIndex}] = numDeviceCUs; + topo.numSubExecutors[{EXE_GPU_ALS, exeIndex}] = numDeviceCUs; topo.closestCpuNumaToGpu[exeIndex] = closestNuma; topo.closestNicsToGpu[exeIndex] = {}; @@ -7362,22 +7627,23 @@ static bool IsConfiguredGid(union ibv_gid const& gid) // NIC Executor int numNics = 0; -#ifdef NIC_EXEC_ENABLED - numNics = GetIbvDeviceList().size(); - for (int exeIndex = 0; exeIndex < numNics; exeIndex++) { - topo.closestCpuNumaToNic[exeIndex] = GetIbvDeviceList()[exeIndex].numaNode; - topo.executorName[{EXE_NIC, exeIndex}] = GetIbvDeviceList()[exeIndex].name; - topo.nicIsActive[exeIndex] = GetIbvDeviceList()[exeIndex].hasActivePort; - if (verbose) { - auto const& nic = GetIbvDeviceList()[exeIndex]; - Log("[INFO] Rank %03d: NIC [%02d/%02d] %s BDF %s NUMA %d active=%s\n", - rank, exeIndex, numNics, nic.name.c_str(), - nic.busId.empty() ? "?" : nic.busId.c_str(), - topo.closestCpuNumaToNic[exeIndex], - nic.hasActivePort ? "yes" : "no"); + if (IsIbvSymbolsReady()) + { + numNics = GetIbvDeviceList().size(); + for (int exeIndex = 0; exeIndex < numNics; exeIndex++) { + topo.closestCpuNumaToNic[exeIndex] = GetIbvDeviceList()[exeIndex].numaNode; + topo.executorName[{EXE_NIC, exeIndex}] = GetIbvDeviceList()[exeIndex].name; + topo.nicIsActive[exeIndex] = GetIbvDeviceList()[exeIndex].hasActivePort; + if (verbose) { + auto const& nic = GetIbvDeviceList()[exeIndex]; + Log("[INFO] Rank %03d: NIC [%02d/%02d] %s BDF %s NUMA %d active=%s\n", + rank, exeIndex, numNics, nic.name.c_str(), + nic.busId.empty() ? "?" : nic.busId.c_str(), + topo.closestCpuNumaToNic[exeIndex], + nic.hasActivePort ? "yes" : "no"); + } } } -#endif topo.numExecutors[EXE_NIC] = topo.numExecutors[EXE_NIC_NEAREST] = numNics; for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { @@ -7402,101 +7668,100 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } // Figure out closest NICs to GPUs -#ifdef NIC_EXEC_ENABLED - // Build up list of NIC bus addresses std::vector ibvAddressList; auto const& ibvDeviceList = GetIbvDeviceList(); - for (auto const& ibvDevice : ibvDeviceList) - ibvAddressList.push_back(ibvDevice.hasActivePort ? ibvDevice.busId : ""); - - // Track how many times a device has been assigned as "closest" - // This allows distributed work across devices using multiple ports (sharing the same busID) - // NOTE: This isn't necessarily optimal, but likely to work in most cases involving multi-port - // Counter example: - // - // G0 prefers (N0,N1), picks N0 - // G1 prefers (N1,N2), picks N1 - // G2 prefers N0, picks N0 - // - // instead of G0->N1, G1->N2, G2->N0 - - std::vector assignedCount(ibvDeviceList.size(), 0); - - // Loop over each GPU to find the closest NIC(s) based on PCIe address - for (int gpuIndex = 0; gpuIndex < numGpus; gpuIndex++) { - if (gpuAddressList[gpuIndex].empty()) continue; - const char* hipPciBusId = gpuAddressList[gpuIndex].c_str(); - - // Find closest NICs - std::set closestNicIdxs = GetNearestDevicesInTree(hipPciBusId, ibvAddressList); - - // Pick the least-used NIC to assign as closest - int closestIdx = -1; - for (auto idx : closestNicIdxs) { - if (closestIdx == -1 || assignedCount[idx] < assignedCount[closestIdx]) - closestIdx = idx; - } - - // The following will only use distance between bus IDs - // to determine the closest NIC to GPU if the PCIe tree approach fails - if (closestIdx < 0) { -#ifdef VERBS_DEBUG - Log("[WARN] Falling back to PCIe bus ID distance to determine proximity\n"); -#endif - int minDistance = std::numeric_limits::max(); - for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { - if (ibvDeviceList[nicIndex].busId != "") { - int distance = GetBusIdDistance(hipPciBusId, ibvDeviceList[nicIndex].busId); - if (distance < minDistance && distance >= 0) { - minDistance = distance; - closestIdx = nicIndex; + if (IsIbvSymbolsReady()) { + for (auto const& ibvDevice : ibvDeviceList) + ibvAddressList.push_back(ibvDevice.hasActivePort ? ibvDevice.busId : ""); + + // Track how many times a device has been assigned as "closest" + // This allows distributed work across devices using multiple ports (sharing the same busID) + // NOTE: This isn't necessarily optimal, but likely to work in most cases involving multi-port + // Counter example: + // + // G0 prefers (N0,N1), picks N0 + // G1 prefers (N1,N2), picks N1 + // G2 prefers N0, picks N0 + // + // instead of G0->N1, G1->N2, G2->N0 + + std::vector assignedCount(ibvDeviceList.size(), 0); + + // Loop over each GPU to find the closest NIC(s) based on PCIe address + for (int gpuIndex = 0; gpuIndex < numGpus; gpuIndex++) { + if (gpuAddressList[gpuIndex].empty()) continue; + const char* hipPciBusId = gpuAddressList[gpuIndex].c_str(); + + // Find closest NICs + std::set closestNicIdxs = GetNearestDevicesInTree(hipPciBusId, ibvAddressList); + + // Pick the least-used NIC to assign as closest + int closestIdx = -1; + for (auto idx : closestNicIdxs) { + if (closestIdx == -1 || assignedCount[idx] < assignedCount[closestIdx]) + closestIdx = idx; + } + + // The following will only use distance between bus IDs + // to determine the closest NIC to GPU if the PCIe tree approach fails + if (closestIdx < 0) { + #ifdef VERBS_DEBUG + Log("[WARN] Falling back to PCIe bus ID distance to determine proximity\n"); + #endif + int minDistance = std::numeric_limits::max(); + for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { + if (ibvDeviceList[nicIndex].busId != "") { + int distance = GetBusIdDistance(hipPciBusId, ibvDeviceList[nicIndex].busId); + if (distance < minDistance && distance >= 0) { + minDistance = distance; + closestIdx = nicIndex; + } } } } + if (closestIdx != -1) { + topo.closestNicsToGpu[gpuIndex].push_back(closestIdx); + assignedCount[closestIdx]++; + } } - if (closestIdx != -1) { - topo.closestNicsToGpu[gpuIndex].push_back(closestIdx); - assignedCount[closestIdx]++; - } - } - // Compute the reverse mapping: closest GPU(s) for each NIC - // Loop over each NIC to find the closest GPU(s) based on PCIe address - for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { - if (!ibvDeviceList[nicIndex].hasActivePort || ibvDeviceList[nicIndex].busId.empty()) { - continue; - } + // Compute the reverse mapping: closest GPU(s) for each NIC + // Loop over each NIC to find the closest GPU(s) based on PCIe address + for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { + if (!ibvDeviceList[nicIndex].hasActivePort || ibvDeviceList[nicIndex].busId.empty()) { + continue; + } - // Find closest GPUs using LCA algorithm - std::set closestGpuIdxs = GetNearestDevicesInTree(ibvDeviceList[nicIndex].busId, gpuAddressList); + // Find closest GPUs using LCA algorithm + std::set closestGpuIdxs = GetNearestDevicesInTree(ibvDeviceList[nicIndex].busId, gpuAddressList); - if (closestGpuIdxs.empty()) { - // Fallback: use bus ID distance - int minDistance = std::numeric_limits::max(); - int closestIdx = -1; + if (closestGpuIdxs.empty()) { + // Fallback: use bus ID distance + int minDistance = std::numeric_limits::max(); + int closestIdx = -1; - for (int gpuIdx = 0; gpuIdx < numGpus; gpuIdx++) { - if (gpuAddressList[gpuIdx].empty()) continue; + for (int gpuIdx = 0; gpuIdx < numGpus; gpuIdx++) { + if (gpuAddressList[gpuIdx].empty()) continue; - int distance = GetBusIdDistance(ibvDeviceList[nicIndex].busId, gpuAddressList[gpuIdx]); - if (distance >= 0 && distance < minDistance) { - minDistance = distance; - closestIdx = gpuIdx; + int distance = GetBusIdDistance(ibvDeviceList[nicIndex].busId, gpuAddressList[gpuIdx]); + if (distance >= 0 && distance < minDistance) { + minDistance = distance; + closestIdx = gpuIdx; + } } - } - if (closestIdx != -1) { - topo.closestGpusToNic[nicIndex].push_back(closestIdx); - } - } else { - // Store all GPUs that are equally close - for (int idx : closestGpuIdxs) { - topo.closestGpusToNic[nicIndex].push_back(idx); + if (closestIdx != -1) { + topo.closestGpusToNic[nicIndex].push_back(closestIdx); + } + } else { + // Store all GPUs that are equally close + for (int idx : closestGpuIdxs) { + topo.closestGpusToNic[nicIndex].push_back(idx); + } } } } -#endif if (verbose) { for (int exeIndex = 0; exeIndex < numGpus; exeIndex++) { @@ -7516,20 +7781,20 @@ static bool IsConfiguredGid(union ibv_gid const& gid) Log("\n"); } } -#ifdef NIC_EXEC_ENABLED - for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { - Log("[INFO] Rank %03d: NIC [%02d/%02d] %s Closest GPUs:", rank, nicIndex, numNics, - ibvDeviceList[nicIndex].name.c_str()); - if (topo.closestGpusToNic[nicIndex].size() == 0) { - Log(" none"); - } else { - for (auto gpuIndex : topo.closestGpusToNic[nicIndex]) { - Log(" %d", gpuIndex); + if (IsIbvSymbolsReady()) { + for (int nicIndex = 0; nicIndex < numNics; nicIndex++) { + Log("[INFO] Rank %03d: NIC [%02d/%02d] %s Closest GPUs:", rank, nicIndex, numNics, + ibvDeviceList[nicIndex].name.c_str()); + if (topo.closestGpusToNic[nicIndex].size() == 0) { + Log(" none"); + } else { + for (auto gpuIndex : topo.closestGpusToNic[nicIndex]) { + Log(" %d", gpuIndex); + } } + Log("\n"); } - Log("\n"); } -#endif } } @@ -7775,6 +8040,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) agent = cpuAgents[exeDevice.exeIndex]; break; case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: + case EXE_GPU_TDM: case EXE_GPU_ALS: if (exeIndex < 0 || exeIndex >= numGpus) return {ERR_FATAL, "GPU index must be between 0 and %d inclusively", numGpus - 1}; agent = gpuAgents[exeIndex]; @@ -8150,6 +8416,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) // Enumerations #undef hipDeviceAttributeClockRate #undef hipDeviceAttributeMultiprocessorCount +#undef hipDeviceAttributeMaxSharedMemoryPerBlock #undef hipDeviceAttributeWarpSize #undef hipErrorPeerAccessAlreadyEnabled #undef hipFuncCachePreferShared @@ -8180,6 +8447,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) #undef hipGetDeviceCount #undef hipGetDeviceProperties #undef hipGetErrorString +#undef hipGetLastError #undef hipHostFree #undef hipHostMalloc #undef hipMalloc diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h new file mode 100644 index 00000000..5175f8c6 --- /dev/null +++ b/src/header/tdmCopy.h @@ -0,0 +1,512 @@ +/* +Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/// \file tdmCopy.h +/// \brief Helper functions to perform HBM to HBM copies via the Tensor Data Mover +/// which copies HBM to a LDS staging buffer, then from LDS out to HBM +/// all without touching cache +/// +/// This is introduced via __device__ level memcpy-like API which can be +/// either blocking or asynchronous, and utilize all warps, or a team of +/// contiguous warps, to allow for other warps to do other tasks +/// +/// \par Quick start +/// \code +/// #include "tdmCopy.h" +/// __shared__ uint8_t staging[N]; // or dynamic extern __shared__ +/// tdm::tdmCopy(dst, src, bytes, staging, N); // block-collective, blocking +/// __syncthreads(); // block-wide visibility +/// \endcode +/// +/// \par Two axes of control +/// - Completion: tdmCopy() (blocking) vs. tdmCopyAsync() + tdmWait() (deferred). +/// - Participation: block-collective (all warps) vs. tdmCopyByTeam() (a +/// contiguous warp range, leaving the other warps free to compute). +/// +/// \par Availability +/// TDM is a hardware feature present on some architectures only. On a target +/// without it, every entry point is `= delete`d: including the header is always +/// fine, but calling any tdm:: function is a hard compile-time error at the call +/// site. See the AVAILABILITY block below. For a runtime/host check (e.g. to pick +/// a kernel before launch), use IsTdmCopySupported(), which is always callable. +/// +/// The design rationale, hardware model, and visibility rules live at the bottom +/// of this file under "IMPLEMENTATION NOTES"; the public API is right here. +#pragma once + +#include +#include +#include + +// ============================================================================ +// AVAILABILITY +// ---------------------------------------------------------------------------- +// TDM is supported on a subset of architectures. Detection is centralized in the +// single macro TDM_SUPPORTED, which is 1 only when we are compiling a device pass +// for a TDM-capable arch AND the compiler exposes the TDM builtin AND the arch's +// descriptor header is on the include path. Keying on __has_builtin / __has_include +// (not just the arch macro) means an older toolchain that predates the builtin, or +// a build without the descriptor header, degrades gracefully instead of failing; +// a new TDM-capable arch works as soon as its target macro is added below. +// +// Each public entry point is individually guarded. When TDM_SUPPORTED is 1, +// the real implementation is compiled. Otherwise every entry point is declared +// `= delete`, so including this header is always fine but CALLING any tdm:: +// function on an unsupported target is a hard compile-time error at the call +// site ("call to deleted function"). The TDM_API / TDM_DELETED macro pair below +// applies that guard uniformly to every declaration. +// ============================================================================ +#ifndef __has_builtin +# define __has_builtin(x) 0 +#endif +#ifndef __has_include +# define __has_include(x) 0 +#endif +// Host-evaluable toolchain capability. TDM_SUPPORTED (below) keys on device arch +// macros (e.g. __gfx1250__) that are never defined during the host pass, so it is +// always 0 in host code and cannot gate the host-side IsTdmCopySupported() check. +// The descriptor header's presence is the reliable host-visible proxy for toolchain +// support (__has_builtin for the amdgcn intrinsic is unreliable in the host/x86 +// pass), and it is also a prerequisite of TDM_SUPPORTED, so it is factored out here. +// Without this host gate, a build whose device pass fell back to the no-op TDM stub +// would still report support on gfx1250 hardware and silently dispatch a no-op copy. +#if __has_include() +# define TDM_TOOLCHAIN_AVAILABLE 1 +#else +# define TDM_TOOLCHAIN_AVAILABLE 0 +#endif + +#if defined(__gfx1250__) && \ + __has_builtin(__builtin_amdgcn_tensor_load_to_lds) && \ + TDM_TOOLCHAIN_AVAILABLE + /* extend: || (defined(__gfxNNNN__) && ...) */ +# define TDM_SUPPORTED 1 +#else +# define TDM_SUPPORTED 0 +#endif + +#if TDM_SUPPORTED +# include // D# descriptor types for the target's TDM +# define TDM_API inline // normal inline declaration (defined below) +# define TDM_DELETED // ... and not deleted +#else +# define TDM_API // no linkage keyword on a deleted declaration +# define TDM_DELETED = delete // unsupported target: any call is a compile error +#endif + +namespace tdm { + +// ============================================================================ +// PUBLIC API +// ============================================================================ + +/// \brief Report whether TDM copies are usable. Always available (never deleted), +/// so it is safe to call on any target as a guard before the copy fns. +/// +/// The answer differs by compilation pass, because TDM availability is a property +/// of the specific GPU arch: +/// - DEVICE code: returns the compile-time constant TDM_SUPPORTED for the arch +/// this device pass was built for. It is a constant expression, so it folds away +/// and can drive `if constexpr` / dead-code elimination of the copy calls. +/// - HOST code: there is no single compile-time answer (a build may target many +/// archs), so it queries the given device's architecture via the HIP runtime and +/// reports whether it is TDM-capable. Returns false if the query fails. +/// +/// \param deviceId HIP device to query (host only; ignored in device code). +/// \return true if tdm:: copies will run on the target/device in question. +/// +/// \code +/// // Host: pick an implementation before launching. +/// if (tdm::IsTdmCopySupported(dev)) launchTdmKernel(...); +/// else launchFallbackKernel(...); +/// \endcode +__host__ __device__ inline bool IsTdmCopySupported(int deviceId = 0) { +#if defined(__HIP_DEVICE_COMPILE__) + (void)deviceId; + return TDM_SUPPORTED; // compile-time constant for this arch pass +#else + hipDeviceProp_t prop; + if (hipGetDeviceProperties(&prop, deviceId) != hipSuccess) return false; + // gcnArchName looks like "gfx1250:sramecc+:xnack-"; match the arch prefix. + // Keep this list in sync with the TDM_SUPPORTED arch condition above. + const char* arch = prop.gcnArchName; + const char* p = "gfx1250"; + while (*p && *arch == *p) { ++arch; ++p; } + // Require BOTH a TDM-capable arch AND a toolchain that actually built TDM + // (otherwise the device pass emitted a no-op stub and enabling the path here + // would silently produce wrong results). + return TDM_TOOLCHAIN_AVAILABLE && (*p == '\0'); +#endif +} + +/// \brief Blocking block-collective copy of [0, sizeBytes): dst <- src. +/// +/// Call from ALL threads of the block with identical arguments (memcpy-style +/// order). Work is partitioned across every warp in the block. On return, the +/// calling warp's TDM ops are complete. +/// +/// \param dst Destination in global memory (HBM). +/// \param src Source in global memory (HBM). +/// \param sizeBytes Number of bytes to copy. +/// \param ldsBuffer Per-block LDS staging area (shared memory). +/// \param ldsBufferBytes Size of \p ldsBuffer in bytes; subdivided among warps. +/// +/// \note For BLOCK-wide visibility, follow with __syncthreads() (and see the +/// visibility notes at the bottom of the file). +/// \warning \p src and \p dst must be GLOBAL (HBM) pointers. Passing a shared / +/// LDS pointer compiles (it decays to void*) but is undefined at run +/// time -- the address is programmed into the descriptor's global field. +__device__ TDM_API void tdmCopy(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes) TDM_DELETED; + +/// \brief Non-blocking variant of tdmCopy(): issue and return. +/// +/// Identical partitioning to tdmCopy(), but does NOT drain on return. The last +/// few TDM ops (bounded by the per-wave queue depth) stay in flight so they +/// overlap with whatever the calling warp does next. Pair with tdmWait(). +/// +/// \see tdmCopy for parameter meanings. +/// \see tdmWait to complete the copy. +__device__ TDM_API void tdmCopyAsync(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes) TDM_DELETED; + +/// \brief Blocking WARP-SPECIALIZED copy performed by one contiguous warp team. +/// +/// Only warps in the half-open range [\p startWarpId, \p stopWarpId) participate; +/// all other warps return immediately and are free to do compute. Work and LDS +/// are partitioned by RANK WITHIN THE TEAM (warpId - startWarpId), so several +/// disjoint teams can each run a different copy concurrently, each with its own +/// dst/src and its own \p ldsBuffer region. +/// +/// \param dst Destination in global memory (HBM). +/// \param src Source in global memory (HBM). +/// \param sizeBytes Number of bytes to copy. +/// \param ldsBuffer THIS team's LDS region (distinct per team). +/// \param ldsBufferBytes Size of this team's LDS region, split among its warps. +/// \param startWarpId First warp of the team (inclusive). +/// \param stopWarpId One past the last warp (clamped to nWarps; ~0u = end). +/// +/// \note This does NOT synchronize the block -- you choose the barrier. A single +/// __syncthreads() after the copy/compute branches is enough for +/// independent compute; use named/arrive-wait barriers for a pipelined +/// producer/consumer so the copy team can run ahead. +/// \warning Give each team (and the compute warps) a NON-OVERLAPPING LDS region; +/// the library trusts the pointer/size you pass. +/// +/// \code +/// const uint32_t warpId = threadIdx.x / warpSize; // (1D block) +/// if (warpId < COPY_WARPS) +/// tdm::tdmCopyByTeam(dst, src, n, staging, teamLdsBytes, 0u, COPY_WARPS); +/// else +/// /* compute -- use LDS past the team's window */; +/// __syncthreads(); +/// \endcode +__device__ TDM_API void tdmCopyByTeam(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) TDM_DELETED; + +/// \brief Non-blocking variant of tdmCopyByTeam(): issue and return. +/// \see tdmCopyByTeam for parameter meanings and team semantics. +/// \see tdmWait to complete the copy (called by each participating warp). +__device__ TDM_API void tdmCopyAsyncByTeam(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) TDM_DELETED; + +/// \brief Drain the CALLING WARP's outstanding TDM ops (TENSORcnt -> 0). +/// +/// TENSORcnt is a per-wave counter, so this waits only on the ops this warp +/// issued -- nothing else. +/// +/// \note Multiple teams need no special handling: each participating warp calls +/// tdmWait() to drain its own ops, and one team's wait has no effect on +/// another's (there is no shared counter). Every ISSUING warp must call it +/// -- a warp cannot drain its teammates' ops. tdmCopy*()/*ByTeam() blocking +/// forms already do this internally. +/// \note For BLOCK-wide visibility follow with __syncthreads(); if copied data is +/// consumed within the block AND the vector head ran, also +/// __threadfence_block() so those ordinary global stores are observed. +__device__ TDM_API void tdmWait() TDM_DELETED; + +} // namespace tdm + +#undef TDM_API +#undef TDM_DELETED + + +// ############################################################################ +// # # +// # IMPLEMENTATION NOTES # +// # # +// ############################################################################ +// +// LAYOUT OF A COPY +// [ head (VECTOR) ][ ---- aligned 256B rows (TDM) ---- ][ tail (TDM 1-D) ] +// * Fixed choices for bandwidth: 4-byte data_size, 256B TDM row width. +// * `head` brings the SOURCE up to a 128B boundary (the direct-copy +// requirement); being the unaligned remainder, it stays a cooperative +// vector copy done by the team's first warp. +// * The aligned bulk is 2D TDM tiles (64 dwords x N rows of 256B). +// * `tail` (< 256B sub-row remainder) is a separate 1-D TDM op at byte +// granularity. TDM's out-of-bounds clamp is per-dimension (rectangular) and +// cannot express "N full rows + a partial row", so the partial row must be +// its own tile rather than riding the bulk descriptor's clamp. +// +// HARDWARE MODEL (why the per-tile waits are REQUIRED) +// * TDM engines: 1 per SIMD-pair -> 2 per WGP; a warp runs on 1 SIMD32 and +// shares its pair's engine. Bandwidth needs >=2 issuing warps/block (both +// engines) and many blocks (many WGPs); a lone warp is latency-bound. +// * Same-wave TDM ops ISSUE in order but their memory effects OVERLAP: up to 3 +// ops are outstanding per wave (that is what TENSORcnt counts). In-order issue +// does NOT serialize completion, so it does NOT make LDS reuse hazard-free. +// * This copy is single-buffered (one LDS window per warp), so each tile is a +// load->store dependency chain on that window: the store reads what the load +// wrote (RAW), and the next tile's load overwrites the window the store is +// still draining (WAR). Both edges need an s_wait_tensorcnt, so issueRows()/ +// issueRow1d() wait after the load and after the store. Consequence: with a +// single window the copy is effectively serialized; tdmCopyAsync() therefore +// overlaps very little today. Regaining overlap needs DOUBLE-BUFFERING (>=2 +// LDS windows per warp) so a load into window B runs while window A stores. +// +// VISIBILITY +// tdmWait() drains only the calling wave's TDM ops. Cross-warp / block-wide +// visibility is the caller's __syncthreads(); the vector head's ordinary +// global stores additionally want __threadfence_block() if consumed within the +// block. Host-after-kernel and grid dependencies are handled by the stream. +// +// BUILTINS: verify names against your tree: +// grep -iE 'tensor_(load|store)|tensorcnt' \ +// /clang/include/clang/Basic/BuiltinsAMDGPU.def +// ============================================================================ + +namespace tdm { + +#if TDM_SUPPORTED // ===== real TDM implementation ================ + +namespace detail { + +constexpr uint32_t WIDTH = 256; // bytes per TDM row (first tile dim) +constexpr uint32_t ELT = 4; // dword +constexpr uint32_t DS4 = 2; // data_size code for 4-byte +constexpr uint32_t DS1 = 0; // data_size code for 1-byte +constexpr uint32_t TD0 = WIDTH / ELT; // 64 elements per row + +// ---- instruction emission (the only arch-specific piece) ------------------- +// The tensor DMA is a single builtin taking the FULL descriptor: five register +// groups plus a constant cache policy. Per the clang reference +// (https://clang.llvm.org/docs/AMDGPUBuiltinReference.html): +// void __builtin_amdgcn_tensor_load_to_lds (v4u32 D0, v8i32 D1, v4i32 D2, +// v4i32 D3, v8i32 D4, int cpol); +// void __builtin_amdgcn_tensor_store_from_lds(); +// D0=GROUP0 (addresses), D1=GROUP1 (2D shape). D2/D3/D4 carry the higher tensor +// dimensions; for a <=2D copy they are simply ZERO vectors ("unused"). This +// mirrors known-good example usage, which passes the group m_bitfields straight +// through (no signed cast) and zero raw vectors for the unused higher dims -- so +// we depend only on GROUP0/GROUP1 existing, not on GROUP2/3/4 by name. +using u32x4 = __attribute__((ext_vector_type(4))) uint32_t; // D0, D2, D3 +using u32x8 = __attribute__((ext_vector_type(8))) uint32_t; // D1, D4 + +__device__ inline void load(const gfx1250_TDM_GROUP0& g0, + const gfx1250_TDM_GROUP1& g1) { + __builtin_amdgcn_tensor_load_to_lds( + g0.m_bitfield, // D0 addresses + g1.m_bitfield, // D1 2D shape + u32x4{}, u32x4{}, u32x8{}, // D2/D3/D4 higher dims: unused (zero) + /*cpol=*/0); +} +__device__ inline void store(const gfx1250_TDM_GROUP0& g0, + const gfx1250_TDM_GROUP1& g1) { + __builtin_amdgcn_tensor_store_from_lds( + g0.m_bitfield, + g1.m_bitfield, + u32x4{}, u32x4{}, u32x8{}, + /*cpol=*/0); +} +__device__ inline void waitTensor0() { __builtin_amdgcn_s_wait_tensorcnt(0); } + +// ---- cooperative vector copy of a small byte range, by one warp ------------ +// All threads of the calling warp participate. Dword-wide where possible; the +// (<4 byte) ragged end is finished by the warp's first thread. +__device__ inline void warpVecCopy(const uint8_t* s, uint8_t* d, size_t n, + uint32_t warpThread, uint32_t warpThreads) { + size_t nd = n >> 2; + const uint32_t* s32 = reinterpret_cast(s); + uint32_t* d32 = reinterpret_cast(d); + for (size_t i = warpThread; i < nd; i += warpThreads) d32[i] = s32[i]; + uint32_t rem = static_cast(n & 3u); + if (rem && warpThread == 0) + for (uint32_t b = 0; b < rem; ++b) d[nd * 4 + b] = s[nd * 4 + b]; +} + +// ---- issue one chunk of whole 256B rows (2D tile) through ONE LDS window. --- +// This staging window is single-buffered, so the two TDM ops form a dependency +// chain that MUST be enforced with TENSORcnt waits -- up to 3 TDM ops are +// outstanding per wave (they overlap), so "same-wave in-order issue" does NOT +// serialize their memory effects: +// * load -> wait: the store reads the LDS the load just wrote (RAW hazard). +// * store -> wait: the caller reuses this same window next iteration; the next +// load must not overwrite LDS the store is still draining (WAR hazard). +__device__ inline void issueRows(uint64_t src, uint32_t lds, uint64_t dst, + uint32_t rows) { + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS4); + g1.tileDim0(TD0); g1.tileDim1(rows); + g1.tensorDim0(TD0); g1.tensorDim1(rows); + g1.tensorDim0Stride(TD0); // rows back-to-back (contiguous) + + // Higher dims (D2/D3/D4) are unused for this 2D tile -> passed as zero inside + // load()/store(), matching known-good example usage. + gfx1250_TDM_GROUP0 g0l(lds, src); + load(g0l, g1); waitTensor0(); // RAW: fill LDS before store reads it + gfx1250_TDM_GROUP0 g0s(lds, dst); + store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse +} + +// ---- issue a sub-row tail (<256B) as a 1-D tile at BYTE granularity. --------- +// Same single-buffered LDS window and the same required RAW/WAR waits as above. +__device__ inline void issueRow1d(uint64_t src, uint32_t lds, uint64_t dst, + uint32_t nbytes) { + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS1); // 1-byte elements: exact length + g1.tileDim0(nbytes); g1.tileDim1(1); + g1.tensorDim0(nbytes); g1.tensorDim1(1); + g1.tensorDim0Stride(nbytes); + + gfx1250_TDM_GROUP0 g0l(lds, src); // unused higher dims -> zero (see load()) + load(g0l, g1); waitTensor0(); // RAW: fill LDS before store reads it + gfx1250_TDM_GROUP0 g0s(lds, dst); + store(g0s, g1); waitTensor0(); // WAR: drain store before window reuse +} + +// ---- core: partition + issue the whole copy for the team [start, stop). ----- +// NO final wait. Work + LDS partition by rank within the team (warpId - start), +// so each team indexes its own `ldsBuffer` from zero. Safe to call collectively +// (warps outside the range return immediately) or only from the team's warps. +__device__ inline void issue(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) { + const uint8_t* s = reinterpret_cast(src); + uint8_t* d = reinterpret_cast(dst); + const uint32_t ldsBase = static_cast(reinterpret_cast(ldsBuffer)); + const uint32_t ldsBytes = static_cast(ldsBufferBytes); // LDS is small + + const uint32_t W = warpSize; + const uint32_t nThreads = blockDim.x * blockDim.y * blockDim.z; + const uint32_t tid = (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + + threadIdx.x; + const uint32_t warpThread = tid % W; // thread index within its warp + const uint32_t warpId = tid / W; + const uint32_t nWarps = (nThreads + W - 1) / W; + + // --- team membership: this warp participates iff in [start, stop) -------- + const uint32_t teamStop = (stopWarpId > nWarps) ? nWarps : stopWarpId; + if (startWarpId >= teamStop || warpId < startWarpId || warpId >= teamStop) + return; // not on this team + const uint32_t rank = warpId - startWarpId; // rank within the team + const uint32_t teamWarps = teamStop - startWarpId; // >= 1 + + // active threads in THIS warp (handles partial final warp); stride for vector. + const uint32_t warpThreads = (nThreads - warpId * W < W) ? (nThreads - warpId * W) : W; + + // --- split the range: [head][ aligned 256B rows ][tail] ------------------ + uint64_t sAddr = reinterpret_cast(s); + uint32_t head = static_cast((128u - (sAddr & 127u)) & 127u); + if (head > sizeBytes) head = static_cast(sizeBytes); + size_t bulk = sizeBytes - head; // starts 128B-aligned + size_t rows = bulk / WIDTH; // whole 256B rows + size_t tdmBytes = rows * WIDTH; + size_t tailOff = head + tdmBytes; + uint32_t tail = static_cast(sizeBytes - tailOff); // < 256B + + // --- edges (team's FIRST warp = rank 0): vector head, TDM tail ----------- + if (rank == 0 && head) warpVecCopy(s, d, head, warpThread, warpThreads); + if (rank == 0 && tail) { + if (ldsBytes >= tail) // stage tail in rank 0's window + issueRow1d(reinterpret_cast(s + tailOff), ldsBase, + reinterpret_cast(d + tailOff), tail); + else + warpVecCopy(s + tailOff, d + tailOff, tail, warpThread, warpThreads); + } + + // --- aligned bulk via TDM ------------------------------------------------ + if (rows == 0) return; // no aligned bulk (edges done) + uint32_t maxByLds = ldsBytes / WIDTH; // #warps we can give a window + if (maxByLds == 0) { // LDS < 256B: vector fallback + if (rank == 0) + warpVecCopy(s + head, d + head, tdmBytes, warpThread, warpThreads); + return; + } + uint32_t issuers = teamWarps < maxByLds ? teamWarps : maxByLds; + uint32_t window = (ldsBytes / issuers) & ~(WIDTH - 1); // per-warp 256B-multiple + uint32_t rowsPerChunk = window / WIDTH; + + if (rank >= issuers) return; // this warp doesn't issue + + // distribute `rows` across issuers by team rank (contiguous row blocks) + size_t base = rows / issuers; + size_t extra = rows % issuers; + size_t myRows = base + (rank < extra ? 1u : 0u); + size_t myStart = rank * base + (rank < extra ? rank : extra); + if (myRows == 0) return; + + uint32_t myLds = ldsBase + rank * window; + uint64_t sBase = reinterpret_cast(s + head) + (uint64_t)myStart * WIDTH; + uint64_t dBase = reinterpret_cast(d + head) + (uint64_t)myStart * WIDTH; + + for (size_t r = 0; r < myRows; r += rowsPerChunk) { + uint32_t chunkRows = (myRows - r < rowsPerChunk) + ? static_cast(myRows - r) : rowsPerChunk; + uint64_t off = (uint64_t)r * WIDTH; + issueRows(sBase + off, myLds, dBase + off, chunkRows); + } +} + +} // namespace detail + +// ---- public API definitions (declared at the top of this file) ------------- + +__device__ inline void tdmWait() { detail::waitTensor0(); } + +__device__ inline void tdmCopyAsync(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes) { + detail::issue(dst, src, sizeBytes, ldsBuffer, ldsBufferBytes, /*start=*/0, /*stop=*/~0u); +} + +__device__ inline void tdmCopy(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes) { + detail::issue(dst, src, sizeBytes, ldsBuffer, ldsBufferBytes, /*start=*/0, /*stop=*/~0u); + tdmWait(); +} + +__device__ inline void tdmCopyAsyncByTeam(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) { + detail::issue(dst, src, sizeBytes, ldsBuffer, ldsBufferBytes, startWarpId, stopWarpId); +} + +__device__ inline void tdmCopyByTeam(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) { + detail::issue(dst, src, sizeBytes, ldsBuffer, ldsBufferBytes, startWarpId, stopWarpId); + tdmWait(); // no-op on any warp that issued nothing / is off-team +} + +#endif // TDM_SUPPORTED +// On an unsupported target the entry points were declared `= delete` at the top, +// so there is nothing to define here -- any call is a compile-time error. + +} // namespace tdm \ No newline at end of file diff --git a/third-party/ibverbs/IbvDynLoad.hpp b/third-party/ibverbs/IbvDynLoad.hpp new file mode 100644 index 00000000..32875278 --- /dev/null +++ b/third-party/ibverbs/IbvDynLoad.hpp @@ -0,0 +1,187 @@ +/* +Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#pragma once + +#include +#include + +#include "IbvHeader.hpp" + +enum IbvLoadStatus { + IBV_OK = 0, + IBV_NO_DMABUF = 1, + IBV_NO_RDMA = 2, +}; + +#define IBV_FN(name, rettype, arglist) rettype(*name)arglist = nullptr; + +namespace { + +IBV_FN(ibv_alloc_pd, ibv_pd*, (ibv_context*)) +IBV_FN(ibv_close_device, int, (ibv_context*)) +IBV_FN(ibv_create_cq, ibv_cq*, (ibv_context*, int, void*, ibv_comp_channel*, int)) +IBV_FN(ibv_create_qp, ibv_qp*, (ibv_pd*, ibv_qp_init_attr*)) +IBV_FN(ibv_dealloc_pd, int, (ibv_pd*)) +IBV_FN(ibv_dereg_mr, int, (ibv_mr*)) +IBV_FN(ibv_destroy_cq, int, (ibv_cq*)) +IBV_FN(ibv_destroy_qp, int, (ibv_qp*)) +IBV_FN(ibv_free_device_list, void, (ibv_device**)) +IBV_FN(ibv_get_device_list, ibv_device**, (int*)) +IBV_FN(ibv_get_device_name, const char*, (ibv_device*)) +IBV_FN(ibv_modify_qp, int, (ibv_qp*, ibv_qp_attr*, int)) +IBV_FN(ibv_open_device, ibv_context*, (ibv_device*)) +IBV_FN(ibv_query_device, int, (ibv_context*, ibv_device_attr*)) +IBV_FN(ibv_query_gid, int, (ibv_context*, uint8_t, int, ibv_gid*)) +IBV_FN(ibv_query_port, int, (ibv_context*, uint8_t, ibv_port_attr*)) +// `ibv_reg_dmabuf_mr` is always declared; whether the underlying symbol +// actually exists in the loaded libibverbs is decided at runtime by tryLoad(). +IBV_FN(ibv_reg_dmabuf_mr, ibv_mr*, (ibv_pd*, uint64_t, size_t, uint64_t, int, int)) +IBV_FN(ibv_reg_mr, ibv_mr*, (ibv_pd*, void*, size_t, int)) +} + +#undef IBV_FN + +struct IbvDynloadState { + std::once_flag once{}; + void* handle = nullptr; + IbvLoadStatus status = IBV_NO_RDMA; + + IbvLoadStatus tryLoad() + { + status = IBV_NO_RDMA; + + handle = dlopen("libibverbs.so.1", RTLD_NOW); + if (handle == nullptr) + return status; + + struct Symbol { void **ppfn; char const *name; }; + + // Core RDMA symbols. Failure of any of these means RDMA is unusable, so we + // tear the whole library back down and report IBV_NO_RDMA. + Symbol coreSymbols[] = { + {(void**)&ibv_alloc_pd, "ibv_alloc_pd"}, + {(void**)&ibv_close_device, "ibv_close_device"}, + {(void**)&ibv_create_cq, "ibv_create_cq"}, + {(void**)&ibv_create_qp, "ibv_create_qp"}, + {(void**)&ibv_dealloc_pd, "ibv_dealloc_pd"}, + {(void**)&ibv_dereg_mr, "ibv_dereg_mr"}, + {(void**)&ibv_destroy_cq, "ibv_destroy_cq"}, + {(void**)&ibv_destroy_qp, "ibv_destroy_qp"}, + {(void**)&ibv_free_device_list, "ibv_free_device_list"}, + {(void**)&ibv_get_device_list, "ibv_get_device_list"}, + {(void**)&ibv_get_device_name, "ibv_get_device_name"}, + {(void**)&ibv_modify_qp, "ibv_modify_qp"}, + {(void**)&ibv_open_device, "ibv_open_device"}, + {(void**)&ibv_query_device, "ibv_query_device"}, + {(void**)&ibv_query_gid, "ibv_query_gid"}, + {(void**)&ibv_query_port, "ibv_query_port"}, + {(void**)&ibv_reg_mr, "ibv_reg_mr"}, + }; + + for (Symbol const& s : coreSymbols) { + void* sym = dlsym(handle, s.name); + if (sym == nullptr) { + // Roll back any pointer already wired so callers don't see a half-loaded library. + for (Symbol const& r : coreSymbols) *r.ppfn = nullptr; + dlclose(handle); + handle = nullptr; + return status; // IBV_NO_RDMA + } + *s.ppfn = sym; + } + + // DMA-BUF probe is independent: missing symbol downgrades to IBV_NO_DMABUF + // but RDMA stays usable. + void* dmabufSym = dlsym(handle, "ibv_reg_dmabuf_mr"); + if (dmabufSym != nullptr) { + *((void**)&ibv_reg_dmabuf_mr) = dmabufSym; + status = IBV_OK; + } else { + ibv_reg_dmabuf_mr = nullptr; + status = IBV_NO_DMABUF; + } + return status; + } +}; + +inline IbvDynloadState& ibvDynloadState() +{ + static IbvDynloadState s; + return s; +} + +inline void IbvEnsureLoaded() +{ + IbvDynloadState& st = ibvDynloadState(); + std::call_once(st.once, [&]() { st.tryLoad(); }); +} + +inline IbvLoadStatus IbvGetLoadStatus() +{ + IbvEnsureLoaded(); + return ibvDynloadState().status; +} + +inline bool IsIbvSymbolsReady() +{ + return IbvGetLoadStatus() != IBV_NO_RDMA; +} + +inline bool IsIbvDmabufPresent() +{ + return IbvGetLoadStatus() == IBV_OK; +} + +inline void* IbvDlHandle() +{ + IbvEnsureLoaded(); + return ibvDynloadState().handle; +} + +inline void IbvUnload() +{ + IbvDynloadState& st = ibvDynloadState(); + if (st.handle != nullptr) { + dlclose(st.handle); + st.handle = nullptr; + st.status = IBV_NO_RDMA; + ibv_alloc_pd = nullptr; + ibv_close_device = nullptr; + ibv_create_cq = nullptr; + ibv_create_qp = nullptr; + ibv_dealloc_pd = nullptr; + ibv_dereg_mr = nullptr; + ibv_destroy_cq = nullptr; + ibv_destroy_qp = nullptr; + ibv_free_device_list = nullptr; + ibv_get_device_list = nullptr; + ibv_get_device_name = nullptr; + ibv_modify_qp = nullptr; + ibv_open_device = nullptr; + ibv_query_device = nullptr; + ibv_query_gid = nullptr; + ibv_query_port = nullptr; + ibv_reg_dmabuf_mr = nullptr; + ibv_reg_mr = nullptr; + } +} diff --git a/third-party/ibverbs/IbvHeader.hpp b/third-party/ibverbs/IbvHeader.hpp new file mode 100644 index 00000000..ca734a5a --- /dev/null +++ b/third-party/ibverbs/IbvHeader.hpp @@ -0,0 +1,586 @@ +/* +* Copyright (c) 2004, 2005 Topspin Communications. All rights reserved. +* Copyright (c) 2004, 2011-2012 Intel Corporation. All rights reserved. +* Copyright (c) 2005, 2006, 2007 Cisco Systems, Inc. All rights reserved. +* Copyright (c) 2005 PathScale, Inc. All rights reserved. +* Copyright (c) 2020 Intel Corporation. All rights reserved. +* Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +* +* This software is available to you under a choice of one of two +* licenses. You may choose to be licensed under the terms of the GNU +* General Public License (GPL) Version 2, available from the file +* COPYING in the main directory of this source tree, or the +* OpenIB.org BSD license below: +* +* Redistribution and use in source and binary forms, with or +* without modification, are permitted provided that the following +* conditions are met: +* +* - Redistributions of source code must retain the above +* copyright notice, this list of conditions and the following +* disclaimer. +* +* - Redistributions in binary form must reproduce the above +* copyright notice, this list of conditions and the following +* disclaimer in the documentation and/or other materials +* provided with the distribution. +* +* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF +* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND +* NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS +* BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN +* ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN +* CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +* SOFTWARE. +*/ + +#pragma once + +#include +#include + +extern "C" { + +// --------------------------------------------------------------------------- +// Opaque handles (forward declarations only) +// --------------------------------------------------------------------------- +struct ibv_pd; +struct ibv_cq; +struct ibv_srq; +struct ibv_ah; +struct ibv_mw; +struct ibv_dm; +struct ibv_xrcd; +struct ibv_comp_channel; + +// --------------------------------------------------------------------------- +// ibv_gid - 16-byte GID. verbs.h declares the inner fields as __be64 (big +// endian). TransferBench only memcpys/broadcasts/compares this opaque blob, +// so plain uint64_t preserves the layout without dragging in . +// --------------------------------------------------------------------------- +union ibv_gid { + uint8_t raw[16]; + struct { + uint64_t subnet_prefix; + uint64_t interface_id; + } global; +}; + +// --------------------------------------------------------------------------- +// Device enumeration types +// --------------------------------------------------------------------------- +enum ibv_node_type { + IBV_NODE_UNKNOWN = -1, + IBV_NODE_CA = 1, + IBV_NODE_SWITCH, + IBV_NODE_ROUTER, + IBV_NODE_RNIC, + IBV_NODE_USNIC, + IBV_NODE_USNIC_UDP, + IBV_NODE_UNSPECIFIED, +}; + +enum ibv_transport_type { + IBV_TRANSPORT_UNKNOWN = -1, + IBV_TRANSPORT_IB = 0, + IBV_TRANSPORT_IWARP, + IBV_TRANSPORT_USNIC, + IBV_TRANSPORT_USNIC_UDP, + IBV_TRANSPORT_UNSPECIFIED, +}; + +enum ibv_atomic_cap { + IBV_ATOMIC_NONE, + IBV_ATOMIC_HCA, + IBV_ATOMIC_GLOB, +}; + +// ibv_device_ops: 2 opaque function pointers; preserved purely for layout. +struct _ibv_device_ops { + struct ibv_context *(*_dummy1)(struct ibv_device *device, int cmd_fd); + void (*_dummy2)(struct ibv_context *context); +}; + +enum { + IBV_SYSFS_NAME_MAX = 64, + IBV_SYSFS_PATH_MAX = 256, +}; + +struct ibv_device { + struct _ibv_device_ops _ops; + enum ibv_node_type node_type; + enum ibv_transport_type transport_type; + char name[IBV_SYSFS_NAME_MAX]; + char dev_name[IBV_SYSFS_NAME_MAX]; + char dev_path[IBV_SYSFS_PATH_MAX]; + char ibdev_path[IBV_SYSFS_PATH_MAX]; +}; + +struct ibv_context_ops { + void *_reserved_pre[11]; // _compat_query_device .. _compat_create_cq + void *poll_cq; // index 11 + void *_reserved_mid[13]; // req_notify_cq .. _compat_destroy_qp + void *post_send; // index 25 +}; + +// We read ->device (offset 0) and dispatch through ->ops. +// The trailing fields(fds, mutex, ...) are intentionally omitted, +// libibverbs allocates and frees these objects. +struct ibv_context { + struct ibv_device *device; + struct ibv_context_ops ops; +}; + +// We only ever read ->context (offset 0) to reach the ops dispatch table. +// The remaining fields are omitted since libibverbs owns the allocation. +struct ibv_cq { + struct ibv_context *context; +}; + +// --------------------------------------------------------------------------- +// Device / port attributes (populated by ibv_query_device / ibv_query_port) +// --------------------------------------------------------------------------- +struct ibv_device_attr { + char fw_ver[64]; + uint64_t node_guid; + uint64_t sys_image_guid; + uint64_t max_mr_size; + uint64_t page_size_cap; + uint32_t vendor_id; + uint32_t vendor_part_id; + uint32_t hw_ver; + int max_qp; + int max_qp_wr; + unsigned int device_cap_flags; + int max_sge; + int max_sge_rd; + int max_cq; + int max_cqe; + int max_mr; + int max_pd; + int max_qp_rd_atom; + int max_ee_rd_atom; + int max_res_rd_atom; + int max_qp_init_rd_atom; + int max_ee_init_rd_atom; + enum ibv_atomic_cap atomic_cap; + int max_ee; + int max_rdd; + int max_mw; + int max_raw_ipv6_qp; + int max_raw_ethy_qp; + int max_mcast_grp; + int max_mcast_qp_attach; + int max_total_mcast_qp_attach; + int max_ah; + int max_fmr; + int max_map_per_fmr; + int max_srq; + int max_srq_wr; + int max_srq_sge; + uint16_t max_pkeys; + uint8_t local_ca_ack_delay; + uint8_t phys_port_cnt; +}; + +enum ibv_mtu { + IBV_MTU_256 = 1, + IBV_MTU_512 = 2, + IBV_MTU_1024 = 3, + IBV_MTU_2048 = 4, + IBV_MTU_4096 = 5, +}; + +enum ibv_port_state { + IBV_PORT_NOP = 0, + IBV_PORT_DOWN = 1, + IBV_PORT_INIT = 2, + IBV_PORT_ARMED = 3, + IBV_PORT_ACTIVE = 4, + IBV_PORT_ACTIVE_DEFER = 5, +}; + +enum { + IBV_LINK_LAYER_UNSPECIFIED, + IBV_LINK_LAYER_INFINIBAND, + IBV_LINK_LAYER_ETHERNET, +}; + +struct ibv_port_attr { + enum ibv_port_state state; + enum ibv_mtu max_mtu; + enum ibv_mtu active_mtu; + int gid_tbl_len; + uint32_t port_cap_flags; + uint32_t max_msg_sz; + uint32_t bad_pkey_cntr; + uint32_t qkey_viol_cntr; + uint16_t pkey_tbl_len; + uint16_t lid; + uint16_t sm_lid; + uint8_t lmc; + uint8_t max_vl_num; + uint8_t sm_sl; + uint8_t subnet_timeout; + uint8_t init_type_reply; + uint8_t active_width; + uint8_t active_speed; + uint8_t phys_state; + uint8_t link_layer; + uint8_t flags; + uint16_t port_cap_flags2; + uint32_t active_speed_ex; +}; + +// --------------------------------------------------------------------------- +// Memory region (populated by ibv_reg_mr / ibv_reg_dmabuf_mr) +// --------------------------------------------------------------------------- +struct ibv_mr { + struct ibv_context *context; + struct ibv_pd *pd; + void *addr; + size_t length; + uint32_t handle; + uint32_t lkey; + uint32_t rkey; +}; + +// --------------------------------------------------------------------------- +// Address handle / global route (used inside ibv_qp_attr.ah_attr) +// --------------------------------------------------------------------------- +struct ibv_global_route { + union ibv_gid dgid; + uint32_t flow_label; + uint8_t sgid_index; + uint8_t hop_limit; + uint8_t traffic_class; +}; + +struct ibv_ah_attr { + struct ibv_global_route grh; + uint16_t dlid; + uint8_t sl; + uint8_t src_path_bits; + uint8_t static_rate; + uint8_t is_global; + uint8_t port_num; +}; + +// --------------------------------------------------------------------------- +// Queue pair init / modify attributes +// --------------------------------------------------------------------------- +enum ibv_qp_type { + IBV_QPT_RC = 2, + IBV_QPT_UC = 3, + IBV_QPT_UD = 4, + IBV_QPT_RAW_PACKET = 8, + IBV_QPT_XRC_SEND = 9, + IBV_QPT_XRC_RECV = 10, + IBV_QPT_DRIVER = 0xff, +}; + +struct ibv_qp_cap { + uint32_t max_send_wr; + uint32_t max_recv_wr; + uint32_t max_send_sge; + uint32_t max_recv_sge; + uint32_t max_inline_data; +}; + +struct ibv_qp_init_attr { + void *qp_context; + struct ibv_cq *send_cq; + struct ibv_cq *recv_cq; + struct ibv_srq *srq; + struct ibv_qp_cap cap; + enum ibv_qp_type qp_type; + int sq_sig_all; +}; + +enum ibv_qp_attr_mask { + IBV_QP_STATE = 1 << 0, + IBV_QP_CUR_STATE = 1 << 1, + IBV_QP_EN_SQD_ASYNC_NOTIFY = 1 << 2, + IBV_QP_ACCESS_FLAGS = 1 << 3, + IBV_QP_PKEY_INDEX = 1 << 4, + IBV_QP_PORT = 1 << 5, + IBV_QP_QKEY = 1 << 6, + IBV_QP_AV = 1 << 7, + IBV_QP_PATH_MTU = 1 << 8, + IBV_QP_TIMEOUT = 1 << 9, + IBV_QP_RETRY_CNT = 1 << 10, + IBV_QP_RNR_RETRY = 1 << 11, + IBV_QP_RQ_PSN = 1 << 12, + IBV_QP_MAX_QP_RD_ATOMIC = 1 << 13, + IBV_QP_ALT_PATH = 1 << 14, + IBV_QP_MIN_RNR_TIMER = 1 << 15, + IBV_QP_SQ_PSN = 1 << 16, + IBV_QP_MAX_DEST_RD_ATOMIC = 1 << 17, + IBV_QP_PATH_MIG_STATE = 1 << 18, + IBV_QP_CAP = 1 << 19, + IBV_QP_DEST_QPN = 1 << 20, + IBV_QP_RATE_LIMIT = 1 << 25, +}; + +enum ibv_qp_state { + IBV_QPS_RESET, + IBV_QPS_INIT, + IBV_QPS_RTR, + IBV_QPS_RTS, + IBV_QPS_SQD, + IBV_QPS_SQE, + IBV_QPS_ERR, + IBV_QPS_UNKNOWN, +}; + +// ibv_qp - layout matches libibverbs through qp_num. TransferBench only ever +// holds ibv_qp* returned by ibv_create_qp and reads qp_num, so the trailing +// libibverbs members (mutex/cond/events_completed) are intentionally omitted: +// we never allocate or sizeof an ibv_qp, and every accessed field sits at its +// real ABI offset. ibv_srq stays opaque (pointer only). +struct ibv_qp { + struct ibv_context *context; + void *qp_context; + struct ibv_pd *pd; + struct ibv_cq *send_cq; + struct ibv_cq *recv_cq; + struct ibv_srq *srq; + uint32_t handle; + uint32_t qp_num; + enum ibv_qp_state state; + enum ibv_qp_type qp_type; +}; + +enum ibv_mig_state { + IBV_MIG_MIGRATED, + IBV_MIG_REARM, + IBV_MIG_ARMED, +}; + +struct ibv_qp_attr { + enum ibv_qp_state qp_state; + enum ibv_qp_state cur_qp_state; + enum ibv_mtu path_mtu; + enum ibv_mig_state path_mig_state; + uint32_t qkey; + uint32_t rq_psn; + uint32_t sq_psn; + uint32_t dest_qp_num; + unsigned int qp_access_flags; + struct ibv_qp_cap cap; + struct ibv_ah_attr ah_attr; + struct ibv_ah_attr alt_ah_attr; + uint16_t pkey_index; + uint16_t alt_pkey_index; + uint8_t en_sqd_async_notify; + uint8_t sq_draining; + uint8_t max_rd_atomic; + uint8_t max_dest_rd_atomic; + uint8_t min_rnr_timer; + uint8_t port_num; + uint8_t timeout; + uint8_t retry_cnt; + uint8_t rnr_retry; + uint8_t alt_port_num; + uint8_t alt_timeout; + uint32_t rate_limit; +}; + +// --------------------------------------------------------------------------- +// Memory access / send flags +// --------------------------------------------------------------------------- +// IBV_ACCESS_RELAXED_ORDERING resolves to IB_UVERBS_ACCESS_OPTIONAL_FIRST, +// which the kernel uAPI defines as (1 << 20). +enum ibv_access_flags { + IBV_ACCESS_LOCAL_WRITE = 1, + IBV_ACCESS_REMOTE_WRITE = (1 << 1), + IBV_ACCESS_REMOTE_READ = (1 << 2), + IBV_ACCESS_REMOTE_ATOMIC = (1 << 3), + IBV_ACCESS_MW_BIND = (1 << 4), + IBV_ACCESS_ZERO_BASED = (1 << 5), + IBV_ACCESS_ON_DEMAND = (1 << 6), + IBV_ACCESS_HUGETLB = (1 << 7), + IBV_ACCESS_FLUSH_GLOBAL = (1 << 8), + IBV_ACCESS_FLUSH_PERSISTENT = (1 << 9), + IBV_ACCESS_RELAXED_ORDERING = (1 << 20), +}; + +enum ibv_wr_opcode { + IBV_WR_RDMA_WRITE, + IBV_WR_RDMA_WRITE_WITH_IMM, + IBV_WR_SEND, + IBV_WR_SEND_WITH_IMM, + IBV_WR_RDMA_READ, + IBV_WR_ATOMIC_CMP_AND_SWP, + IBV_WR_ATOMIC_FETCH_AND_ADD, + IBV_WR_LOCAL_INV, + IBV_WR_BIND_MW, + IBV_WR_SEND_WITH_INV, + IBV_WR_TSO, + IBV_WR_DRIVER1, + IBV_WR_FLUSH = 14, + IBV_WR_ATOMIC_WRITE = 15, +}; + +enum ibv_send_flags { + IBV_SEND_FENCE = 1 << 0, + IBV_SEND_SIGNALED = 1 << 1, + IBV_SEND_SOLICITED = 1 << 2, + IBV_SEND_INLINE = 1 << 3, + IBV_SEND_IP_CSUM = 1 << 4, +}; + +// --------------------------------------------------------------------------- +// Scatter/gather and work request (consumed by ibv_post_send) +// --------------------------------------------------------------------------- +struct ibv_sge { + uint64_t addr; + uint32_t length; + uint32_t lkey; +}; + +// Forward decl needed by ibv_send_wr.bind_mw (kept for layout). Mirrors +// verbs.h's struct ibv_mw_bind_info exactly. +struct ibv_mw_bind_info { + struct ibv_mr *mr; + uint64_t addr; + uint64_t length; + unsigned int mw_access_flags; +}; + +// Full ABI-exact ibv_send_wr. TransferBench only sets the rdma arm of `wr`, +// but the union's overall size must match the system layout because the +// driver may write through the entire struct. +struct ibv_send_wr { + uint64_t wr_id; + struct ibv_send_wr *next; + struct ibv_sge *sg_list; + int num_sge; + enum ibv_wr_opcode opcode; + unsigned int send_flags; + union { + uint32_t imm_data; + uint32_t invalidate_rkey; + }; + union { + struct { + uint64_t remote_addr; + uint32_t rkey; + } rdma; + struct { + uint64_t remote_addr; + uint64_t compare_add; + uint64_t swap; + uint32_t rkey; + } atomic; + struct { + struct ibv_ah *ah; + uint32_t remote_qpn; + uint32_t remote_qkey; + } ud; + } wr; + union { + struct { + uint32_t remote_srqn; + } xrc; + } qp_type; + union { + struct { + struct ibv_mw *mw; + uint32_t rkey; + struct ibv_mw_bind_info bind_info; + } bind_mw; + struct { + void *hdr; + uint16_t hdr_sz; + uint16_t mss; + } tso; + }; +}; + +// --------------------------------------------------------------------------- +// Completion queue entry (populated by ibv_poll_cq) +// --------------------------------------------------------------------------- +enum ibv_wc_status { + IBV_WC_SUCCESS, + IBV_WC_LOC_LEN_ERR, + IBV_WC_LOC_QP_OP_ERR, + IBV_WC_LOC_EEC_OP_ERR, + IBV_WC_LOC_PROT_ERR, + IBV_WC_WR_FLUSH_ERR, + IBV_WC_MW_BIND_ERR, + IBV_WC_BAD_RESP_ERR, + IBV_WC_LOC_ACCESS_ERR, + IBV_WC_REM_INV_REQ_ERR, + IBV_WC_REM_ACCESS_ERR, + IBV_WC_REM_OP_ERR, + IBV_WC_RETRY_EXC_ERR, + IBV_WC_RNR_RETRY_EXC_ERR, + IBV_WC_LOC_RDD_VIOL_ERR, + IBV_WC_REM_INV_RD_REQ_ERR, + IBV_WC_REM_ABORT_ERR, + IBV_WC_INV_EECN_ERR, + IBV_WC_INV_EEC_STATE_ERR, + IBV_WC_FATAL_ERR, + IBV_WC_RESP_TIMEOUT_ERR, + IBV_WC_GENERAL_ERR, + IBV_WC_TM_ERR, + IBV_WC_TM_RNDV_INCOMPLETE, +}; + +enum ibv_wc_opcode { + IBV_WC_SEND, + IBV_WC_RDMA_WRITE, + IBV_WC_RDMA_READ, + IBV_WC_COMP_SWAP, + IBV_WC_FETCH_ADD, + IBV_WC_BIND_MW, + IBV_WC_LOCAL_INV, + IBV_WC_TSO, + IBV_WC_FLUSH, + IBV_WC_ATOMIC_WRITE = 9, + IBV_WC_RECV = 1 << 7, + IBV_WC_RECV_RDMA_WITH_IMM, + IBV_WC_TM_ADD, + IBV_WC_TM_DEL, + IBV_WC_TM_SYNC, + IBV_WC_TM_RECV, + IBV_WC_TM_NO_TAG, + IBV_WC_DRIVER1, + IBV_WC_DRIVER2, + IBV_WC_DRIVER3, +}; + +struct ibv_wc { + uint64_t wr_id; + enum ibv_wc_status status; + enum ibv_wc_opcode opcode; + uint32_t vendor_err; + uint32_t byte_len; + union { + uint32_t imm_data; + uint32_t invalidated_rkey; + }; + uint32_t qp_num; + uint32_t src_qp; + unsigned int wc_flags; + uint16_t pkey_index; + uint16_t slid; + uint8_t sl; + uint8_t dlid_path_bits; +}; + +static inline int ibv_poll_cq(struct ibv_cq *cq, int num_entries, struct ibv_wc *wc) +{ + return ((int (*)(struct ibv_cq*, int, struct ibv_wc*))cq->context->ops.poll_cq)( + cq, num_entries, wc); +} + +static inline int ibv_post_send(struct ibv_qp *qp, struct ibv_send_wr *wr, + struct ibv_send_wr **bad_wr) +{ + return ((int (*)(struct ibv_qp*, struct ibv_send_wr*, struct ibv_send_wr**)) + qp->context->ops.post_send)(qp, wr, bad_wr); +} +} // extern "C"