From 85bd7ffc082e2e5733e4ae507cbfef47f280374e Mon Sep 17 00:00:00 2001 From: alex-breslow-amd Date: Fri, 12 Jun 2026 15:05:24 -0700 Subject: [PATCH 1/5] Add Logic for Using TDM and Async Loads and Stores (#322) * Put in stubs * Add kernel and fix up * Finalize tunability of threads per block and number of blocks. * Implement pipelining * Add debugging logic * Fix segfault * Disable debug prints * Add missing timing information * Add TDM_BLOCK_SIZE, TDM_MAX_LDS_BYTES, and TDM_PIPELINED env vars to control TDM kernel execution * Not functioning but proof of concept async load/store implementation * Fix async load store kernel MVP * Potential fix for pull request finding Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Tim <43156029+AtlantaPepsi@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --- Makefile | 8 + src/client/EnvVars.hpp | 23 ++ src/client/Presets/Help.hpp | 24 +- src/client/Utilities.hpp | 2 + src/header/TransferBench.hpp | 419 ++++++++++++++++++++++++++++++++++- src/header/tdm.h | 355 +++++++++++++++++++++++++++++ 6 files changed, 815 insertions(+), 16 deletions(-) create mode 100644 src/header/tdm.h diff --git a/Makefile b/Makefile index db40ed20..46858abc 100644 --- a/Makefile +++ b/Makefile @@ -6,7 +6,15 @@ 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 diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 34940b69..69bd23be 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -124,6 +124,11 @@ 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 tdmPipelined; // Use the pipelined (depth=2) TDM tensor kernel (0=non-pipelined, 1=pipelined) + // Developer features int gpuMaxHwQueues; // Tracks GPU_MAX_HW_QUEUES environment variable @@ -173,6 +178,9 @@ 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); + tdmPipelined = GetEnvVar("TDM_PIPELINED" , 0); useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); @@ -394,6 +402,9 @@ class EnvVars 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(" 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_PIPELINED - 1=use pipelined (depth=2) TDM kernel, 0=use non-pipelined kernel\n"); printf(" USE_HIP_EVENTS - Use HIP events for GFX executor timing\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"); @@ -541,6 +552,14 @@ class EnvVars "%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_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_PIPELINED", tdmPipelined, + "Using %s TDM tensor kernel", tdmPipelined ? "pipelined (depth=2)" : "non-pipelined"); Print("USE_HIP_EVENTS", useHipEvents, "Using %s for GFX/DMA Executor timing", useHipEvents ? "HIP events" : "CPU wall time"); Print("USE_HSA_DMA", useHsaDma, @@ -749,6 +768,10 @@ class EnvVars cfg.nic.trafficClass = nicTrafficClass; cfg.nic.roceVersion = roceVersion; + cfg.tdm.blockSize = tdmBlockSize; + cfg.tdm.maxLDSBytes = tdmMaxLdsBytes; + cfg.tdm.pipelined = tdmPipelined; + return cfg; } }; diff --git a/src/client/Presets/Help.hpp b/src/client/Presets/Help.hpp index 26ede846..de92ee9e 100644 --- a/src/client/Presets/Help.hpp +++ b/src/client/Presets/Help.hpp @@ -37,13 +37,15 @@ 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("# Seven 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("# 6) TDM GPU threadblock/Compute Unit (CU)\n"); + printf("# 7) 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 +73,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 kernel (Indexed from 0 to # GPUs - 1)\n"); + printf("# - L: 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 +112,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 tensor-op 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->L0->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/Utilities.hpp b/src/client/Utilities.hpp index 9730ace6..77a44b3c 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_ASYNC_TENSOR: return "AT"; // async tensor kernel path + case EXE_GPU_ASYNC_MEMOPS: return "AL"; // async load/store kernel path default: return "N/A"; } } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 5f22f52c..d5072b1b 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -25,6 +25,10 @@ THE SOFTWARE. #include #include #include +#include +#include +#include +#include #include #include #include @@ -77,9 +81,9 @@ THE SOFTWARE. #ifdef AMD_SMI_ENABLED #include "amd_smi/amdsmi.h" #endif +#include "tdm.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 +111,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_ASYNC_TENSOR = 6, ///< GPU async kernel — tensor-op path (subExecutor — reserved; stub uses 1-D launch) + EXE_GPU_ASYNC_MEMOPS = 7, ///< GPU async kernel — load/store path (subExecutor — reserved; stub uses 1-D launch) }; - char const ExeTypeStr[7] = "CGDINB"; + char const ExeTypeStr[9] = "CGDINBTL"; 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_ASYNC_TENSOR || e == EXE_GPU_ASYNC_MEMOPS; + } inline bool IsNicExeType(ExeType e){ return e == EXE_NIC || e == EXE_NIC_NEAREST; } /** @@ -275,6 +284,13 @@ 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 64) + int maxLDSBytes = INT_MAX; ///< Maximum number of bytes of __shared__ memory to use + int pipelined = 0; ///< Whether to call the pipelined TDM kernel + }; + /** * Configuration options for performing Transfers @@ -287,6 +303,7 @@ namespace TransferBench GfxOptions gfx; ///< GFX executor options DmaOptions dma; ///< DMA executor options NicOptions nic; ///< NIC executor options + TdmOptions tdm; ///< TDM executor options }; /** @@ -2242,7 +2259,8 @@ 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 (t.exeDevice.exeType == EXE_GPU_GFX || t.exeDevice.exeType == EXE_CPU || t.exeDevice.exeType == EXE_GPU_BDMA || + t.exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || t.exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { 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 +2318,27 @@ namespace { hasFatalError = true; } break; + case EXE_GPU_ASYNC_TENSOR: + case EXE_GPU_ASYNC_MEMOPS: + if (t.srcs.size() != 1 || t.dsts.size() != 1) { + errors.push_back({ERR_FATAL, + "Transfer %d: Async GPU kernel executor 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: Async GPU 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: Async GPU kernel executor does not support subindices yet", i}); + hasFatalError = true; + } + break; case EXE_GPU_GFX: if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { errors.push_back({ERR_FATAL, @@ -2671,6 +2710,22 @@ namespace { } break; } + case EXE_GPU_ASYNC_TENSOR: + case EXE_GPU_ASYNC_MEMOPS: + { + int numGpuSubExec = GetNumSubExecutors(exeDevice); + if (totalSubExecs[exeDevice] > numGpuSubExec) + errors.push_back({ERR_WARN, + "GPU %d (async kernel) 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 (async kernel) 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 @@ -4327,13 +4382,16 @@ 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) + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || + exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) && 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_ASYNC_TENSOR || + exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) ? exeInfo.resources.size() : 1; exeInfo.streams.resize(numStreamsToUse); @@ -4351,7 +4409,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { + if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents || exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || + exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { exeInfo.startEvents.resize(numStreamsToUse); exeInfo.stopEvents.resize(numStreamsToUse); for (int i = 0; i < numStreamsToUse; ++i) { @@ -4557,11 +4616,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } // Teardown additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || + exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) && exeDevice.exeRank == localRank) { for (auto stream : exeInfo.streams) ERR_CHECK(hipStreamDestroy(stream)); - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { + if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents || exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || + exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { for (auto event : exeInfo.startEvents) ERR_CHECK(hipEventDestroy(event)); for (auto event : exeInfo.stopEvents) @@ -4902,6 +4963,203 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } } + //---------------------------------------------------------------------------- +#if defined(__NVCC__) + __global__ void GpuAsyncTensorOpsStubKernel(float const* __restrict__ src, + float* __restrict__ dst, + size_t numFloats) + { + size_t const gid = blockIdx.x * blockDim.x + threadIdx.x; + size_t const stride = gridDim.x * blockDim.x; + for (size_t i = gid; i < numFloats; i += stride) + dst[i] = src[i]; + } +#else +// The TDM API does not have a function to set the transfer size, so we need to do it manually. +__device__ void SetTransferSize(gfx1250_TDM_GROUP1& group1, int numElements){ + group1.tensorDim0(numElements); + group1.tensorDim0Stride(numElements); + group1.tileDim0(numElements); +} +namespace { + __constant__ constexpr bool verbose = false; // Turn off to disable debug prints from kernel TDM kernels. +} + +// numElementsPerTile is the number of elements to process per tile. Its maximum value is the shared memory size / number of waves per workgroup. +__global__ void GpuAsyncTensorOpsKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerTile){ + extern __shared__ __align__(128) float shmem[]; + int waveId = threadIdx.x / warpSize; + int numWavesPerBlock = blockDim.x / warpSize; + size_t itemsProcessedPerGridIteration = numElementsPerTile * numWavesPerBlock * gridDim.x; + + float* shmemPtr = static_cast(shmem) + numElementsPerTile * waveId; + // Local per-wave source and destination pointers + const float* srcPtr = src + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + float* dstPtr = dst + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + gfx1250_TDM_GROUP0 group0; + + group0.ldsAddr((uintptr_t)shmemPtr); + + gfx1250_TDM_GROUP1 group1; + group1.dataSize(2); // Log2 of the element size in bytes, so 2 for float + SetTransferSize(group1, numElementsPerTile); + + constexpr __hip_uint32x4 empty_x4{}; + constexpr __hip_uint32x8 empty_x8{}; + size_t elementsToProcess = numElementsPerTile; + while(srcPtr < src + numElements){ + // Handle the last tile of the block, which may be less than num_elements_per_tile. + if(src + numElements - srcPtr < numElementsPerTile){ + size_t remainingElements = src + numElements - srcPtr; + SetTransferSize(group1, remainingElements); + if constexpr(verbose) elementsToProcess = remainingElements; + } + // Copy from global memory to LDS + group0.globalAddr((uintptr_t)srcPtr); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Block %d, Wave %d: Loading %zu elements from %p to shared memory at %p\n", blockIdx.x, waveId, elementsToProcess, srcPtr, shmemPtr); + __builtin_amdgcn_tensor_load_to_lds(group0.m_bitfield, group1.m_bitfield, empty_x4, empty_x4, empty_x8, 0); + __builtin_amdgcn_s_wait_tensorcnt(0); + + // write back from LDS to global + group0.globalAddr((uintptr_t)dstPtr); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Block %d, Wave %d: Storing %zu elements from shared memory at %p to %p\n", blockIdx.x, waveId, elementsToProcess, shmemPtr, dstPtr); + __builtin_amdgcn_tensor_store_from_lds(group0.m_bitfield, group1.m_bitfield, empty_x4, empty_x4, empty_x8, 0); + __builtin_amdgcn_s_wait_tensorcnt(0); + srcPtr += itemsProcessedPerGridIteration; + dstPtr += itemsProcessedPerGridIteration; + } +} + +struct TileMover { + gfx1250_TDM_GROUP0 group0; + gfx1250_TDM_GROUP1 group1; + float* dstPtr{nullptr}; + float* shmemPtr{nullptr}; // TODO: Remove this + int numElementsToProcess{0}; // TODO: Remove this + + __device__ void LoadTile(float* shmemPtr, const float* srcPtr, float* dstPtr, int numElementsToProcess) { + group0.ldsAddr((uintptr_t)shmemPtr); + group0.globalAddr((uintptr_t)srcPtr); + group1.dataSize(2); + SetTransferSize(group1, numElementsToProcess); + this -> dstPtr = dstPtr; this -> shmemPtr = shmemPtr; this -> numElementsToProcess = numElementsToProcess; + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Loading %d elements from %p to %p\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess, srcPtr, shmemPtr); + __builtin_amdgcn_tensor_load_to_lds(group0.m_bitfield, group1.m_bitfield, __hip_uint32x4{}, __hip_uint32x4{}, __hip_uint32x8{}, 0); + } + __device__ void StoreTile() { + assert(dstPtr != nullptr); + group0.globalAddr((uintptr_t)dstPtr); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Storing %d elements from %p to %p\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess, shmemPtr, dstPtr); + __builtin_amdgcn_tensor_store_from_lds(group0.m_bitfield, group1.m_bitfield, __hip_uint32x4{}, __hip_uint32x4{}, __hip_uint32x8{}, 0); + } +}; + +// numElementsPerTile is the number of elements to process per tile. Its maximum value is the shared memory size / number of waves per workgroup. +__global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerSubtile){ + extern __shared__ __align__(128) float shmem[]; + int waveId = threadIdx.x / warpSize; + int numWavesPerBlock = blockDim.x / warpSize; // This only works because the blockDim.x is fixed. + constexpr int pipelineDepth = 2; + + int numElementsPerTile = pipelineDepth * numElementsPerSubtile; + size_t itemsProcessedPerGridIteration = numElementsPerTile * numWavesPerBlock * gridDim.x; + + float* shmemPtr = static_cast(shmem) + numElementsPerTile * waveId; + // Local per-wave source and destination pointers + const float* srcPtr = src + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + float* dstPtr = dst + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); + + // Skip waves whose initial position is already past the end of the buffer. + if (srcPtr >= src + numElements){ + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Skipping wave %d\n", threadIdx.x / warpSize, blockIdx.x, waveId); + return; + } + + int numElementsToProcess = std::min(numElementsPerSubtile, (int)(src + numElements - srcPtr)); + TileMover tileMovers[pipelineDepth]{}; + unsigned int slot = 0; + tileMovers[slot].LoadTile(shmemPtr, srcPtr, dstPtr, numElementsToProcess); + while(srcPtr + (slot ^ 1) * numElementsPerSubtile < src + numElements){ + slot ^= 1; + // Handle the last tile of the block, which may be less than num_elements_per_tile. + if(src + numElements - (srcPtr + slot * numElementsPerSubtile) < numElementsPerSubtile){ + numElementsToProcess = src + numElements - (srcPtr + slot * numElementsPerSubtile); + if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Small Subtile: Warp %d Block %d: Num Elements To Process: %d\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess); + } + // Copy from global memory to LDS + tileMovers[slot].LoadTile(shmemPtr + slot * numElementsPerSubtile, srcPtr + slot * numElementsPerSubtile, dstPtr + slot * numElementsPerSubtile, numElementsToProcess); + __builtin_amdgcn_s_wait_tensorcnt(1); + + // write back from LDS to global for the other slot + tileMovers[slot^1].StoreTile(); + __builtin_amdgcn_s_wait_tensorcnt(1); + srcPtr += itemsProcessedPerGridIteration * slot; + dstPtr += itemsProcessedPerGridIteration * slot; + } + __builtin_amdgcn_s_wait_tensorcnt(0); + tileMovers[slot].StoreTile(); + __builtin_amdgcn_s_wait_tensorcnt(0); +} +#endif + + // Source for definitions: https://github.com/llvm/llvm-project/blob/4e3bac3ea2cc6fd778d53e317dd9fc27c1ddfc4f/clang/include/clang/Basic/BuiltinsAMDGPU.td#L956 + // and internal ISA documentation + using uint32x4 = __attribute__((__vector_size__(4 * sizeof(int)))) int; + // Loads 16 bytes/lane from global (src) into LDS (dst). + __device__ void asyncLoadX4(float const* src, float* dst){ + __builtin_amdgcn_global_load_async_to_lds_b128( + (__attribute__((address_space(1))) uint32x4*)src, + (__attribute__((address_space(3))) uint32x4*)dst, 0, 0); + } + // Stores 16 bytes/lane from LDS (src) to global (dst). + __device__ void asyncStoreX4(float const* src, float* dst){ + __builtin_amdgcn_global_store_async_from_lds_b128( + (__attribute__((address_space(1))) uint32x4*)dst, + (__attribute__((address_space(3))) uint32x4*)src, 0, 0); + } + + // numElementsPerTile should be a multiple of 128, given the asyncLoadX4 and asyncStoreX4 functions operate on a 32-lane warp and 4 elements per lane (512 bytes total) + __global__ void GpuAsyncLoadStoreKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerTile) + { + extern __shared__ __align__(128) float shmem[]; + int waveId = threadIdx.x / warpSize; + int laneId = threadIdx.x % warpSize; + int numWavesPerBlock = blockDim.x / warpSize; + size_t itemsProcessedPerGridIteration = (size_t)numElementsPerTile * numWavesPerBlock * gridDim.x; + constexpr size_t stride_per_lane = sizeof(float4) / sizeof(float); // 4 elements per lane + // Number of elements a full 32-lane warp moves per b128 async op (32 * 4 = 128). + size_t elements_per_dwordx4_load_store = warpSize * stride_per_lane; + + float* shmemPtr = shmem + numElementsPerTile * waveId + laneId * stride_per_lane; + // Wave-tile base (lane-independent) used to gate the loop and compute the partial-tile size. + size_t tileBase = (size_t)numElementsPerTile * (waveId + (size_t)blockIdx.x * numWavesPerBlock); + + while(tileBase < numElements){ + // Handle the last tile, which may be less than numElementsPerTile. + size_t elementsToProcess = numElementsPerTile; + if(numElements - tileBase < (size_t)numElementsPerTile) + elementsToProcess = numElements - tileBase; + + const float* srcPtr = src + tileBase + laneId * stride_per_lane; + float* dstPtr = dst + tileBase + laneId * stride_per_lane; + + // Copy from global memory to LDS + for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store) { + asyncLoadX4(srcPtr + i, shmemPtr + i); + __builtin_amdgcn_s_wait_asynccnt(0); + } + + // Write back from LDS to global + for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store){ + asyncStoreX4(shmemPtr + i, dstPtr + i); + __builtin_amdgcn_s_wait_asynccnt(0); + } + + tileBase += itemsProcessedPerGridIteration; + } +} + + // Simplified Kernel for GFX execution for copies only template __global__ void __launch_bounds__(LAUNCH_BOUND) @@ -5489,6 +5747,132 @@ static bool IsConfiguredGid(union ibv_gid const& gid) return ERR_NONE; } + static ErrResult ExecuteGpuAsyncKernelTransfer(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + hipStream_t const stream, + hipEvent_t const startEvent, + hipEvent_t const stopEvent, + bool const tensorFlavor, + int const numSubExecs, + TransferResources& rss) + { + ERR_CHECK(hipSetDevice(exeIndex)); + int const initOffset = cfg.data.byteOffset / sizeof(float); + size_t const numFloats = rss.numBytes / sizeof(float); + float const* __restrict__ src = rss.srcMem[0] + initOffset; + float* __restrict__ dst = rss.dstMem[0] + initOffset; + + int const threads = tensorFlavor ? cfg.tdm.blockSize : cfg.gfx.blockSize; + int const blocks = numSubExecs; + + auto cpuStart = std::chrono::high_resolution_clock::now(); + + int subIterations = 0; + do { + if (startEvent) + ERR_CHECK(hipEventRecord(startEvent, stream)); +#if defined(__NVCC__) + GpuAsyncTensorOpsStubKernel<<>>(src, dst, numFloats); +#else + bool usePipelinedTensorOps = (cfg.tdm.pipelined != 0); + auto gpuKernel = tensorFlavor ? (usePipelinedTensorOps ? GpuAsyncPipelinedTensorOpsKernel : GpuAsyncTensorOpsKernel) : GpuAsyncLoadStoreKernel; + if (rss.numBytes % sizeof(float) != 0) + return {ERR_FATAL, "Async tensor executor (TDM): numBytes (%zu) must be a multiple of %zu", rss.numBytes, sizeof(float)}; + + + hipDeviceProp_t props{}; + ERR_CHECK(hipGetDeviceProperties(&props, exeIndex)); + int const warpSize = props.warpSize; + int maxShmem = 0; + ERR_CHECK(hipDeviceGetAttribute(&maxShmem, hipDeviceAttributeMaxSharedMemoryPerBlock, exeIndex)); + maxShmem = std::min(maxShmem, cfg.tdm.maxLDSBytes); + int kWavesPerWorkgroup = threads / warpSize; + int pipelineDepth = usePipelinedTensorOps ? 2 : 1; + int numFloatsPerTile = std::max(1, maxShmem / (sizeof(float) * kWavesPerWorkgroup * pipelineDepth)); + unsigned int shmemBytes = numFloatsPerTile * kWavesPerWorkgroup * sizeof(float) * pipelineDepth; + hipLaunchKernelGGL(gpuKernel, dim3(blocks), dim3(threads), + shmemBytes, stream, src, dst, numFloats, + numFloatsPerTile); +#endif + ERR_CHECK(hipGetLastError()); + if (stopEvent) + ERR_CHECK(hipEventRecord(stopEvent, stream)); + ERR_CHECK(hipStreamSynchronize(stream)); + } while (++subIterations != cfg.general.numSubIterations); + + if (iteration >= 0) { + double deltaMsec; + if (startEvent && stopEvent) { + float gpuDeltaMsec; + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); + deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations; + } else { + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 + / cfg.general.numSubIterations; + } + rss.totalDurationMsec += deltaMsec; + if (cfg.general.recordPerIteration) + rss.perIterMsec.push_back(deltaMsec); + } + return ERR_NONE; + } + + static ErrResult RunGpuAsyncKernelExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + bool const tensorFlavor, + 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 = (!exeInfo.startEvents.empty()) ? exeInfo.startEvents[i] : nullptr; + hipEvent_t stopEv = (!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, + ExecuteGpuAsyncKernelTransfer, + iteration, + std::cref(cfg), + exeIndex, + exeInfo.streams[i], + startEv, + stopEv, + tensorFlavor, + 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; + } + + static ErrResult RunGpuAsyncTensorExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + ExeInfo& exeInfo) + { + return RunGpuAsyncKernelExecutor(iteration, cfg, exeIndex, true, exeInfo); + } + + static ErrResult RunGpuAsyncMemOpsExecutor(int const iteration, + ConfigOptions const& cfg, + int const exeIndex, + ExeInfo& exeInfo) + { + return RunGpuAsyncKernelExecutor(iteration, cfg, exeIndex, false, exeInfo); + } + // DMA Executor-related functions //======================================================================================== @@ -5713,6 +6097,8 @@ 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_ASYNC_TENSOR: return RunGpuAsyncTensorExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_ASYNC_MEMOPS: return RunGpuAsyncMemOpsExecutor(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); @@ -6413,6 +6799,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_ASYNC_TENSOR: + case EXE_GPU_ASYNC_MEMOPS: + 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 @@ -7300,6 +7692,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid) topo.numExecutors[EXE_GPU_GFX] = numGpus; topo.numExecutors[EXE_GPU_DMA] = numGpus; topo.numExecutors[EXE_GPU_BDMA] = numGpus; + topo.numExecutors[EXE_GPU_ASYNC_TENSOR] = numGpus; + topo.numExecutors[EXE_GPU_ASYNC_MEMOPS] = numGpus; std::vector gpuArchNames(numGpus); @@ -7323,6 +7717,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_ASYNC_TENSOR, exeIndex}] = gpuName + " [async tensor]"; + topo.executorName[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = gpuName + " [async LD/ST]"; #if !defined(__NVCC__) hsa_agent_t gpuAgent = gpuAgents[exeIndex]; @@ -7352,9 +7748,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_ASYNC_TENSOR, exeIndex}] = 0; + topo.numExecutorSubIndices[{EXE_GPU_ASYNC_MEMOPS, 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_ASYNC_TENSOR, exeIndex}] = numDeviceCUs; + topo.numSubExecutors[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = numDeviceCUs; topo.closestCpuNumaToGpu[exeIndex] = closestNuma; topo.closestNicsToGpu[exeIndex] = {}; @@ -7775,6 +8175,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_ASYNC_TENSOR: case EXE_GPU_ASYNC_MEMOPS: if (exeIndex < 0 || exeIndex >= numGpus) return {ERR_FATAL, "GPU index must be between 0 and %d inclusively", numGpus - 1}; agent = gpuAgents[exeIndex]; diff --git a/src/header/tdm.h b/src/header/tdm.h new file mode 100644 index 00000000..c55a2bdd --- /dev/null +++ b/src/header/tdm.h @@ -0,0 +1,355 @@ +/* +Copyright (c) 2025 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. +*/ + +#ifndef HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H +#define HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H + +#if !defined(__HIPCC_RTC__) +/* hip_runtime_api.h defines hipMemoryType before including driver_types.h; do not include + * driver_types.h alone or HIP 7.x fails with unknown type hipMemoryType. */ +#include +#include +#if __cplusplus +#include +#else +#include +#endif +#endif + +using __hip_uint32x4 = __NATIVE_VECTOR__(4, uint32_t); +using __hip_uint32x8 = __NATIVE_VECTOR__(8, uint32_t); + +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wshift-count-overflow" // we're going to be using overflow on purpose + +union gfx1250_TDM_GROUP0 +{ + __device__ constexpr gfx1250_TDM_GROUP0() : m_bitfield{0x1, 0, 0, 0x80000000} + {} + + __device__ gfx1250_TDM_GROUP0( uintptr_t lds_addr_in, + uintptr_t global_addr_in, + uint32_t gather_idx_size_in = 0, + uint32_t gather_mode_in = 0 ) + { + globalAddr(global_addr_in); + m_count = 1; + m_lds_addr = lds_addr_in; + m_type = 2; + m_gather_index_size = gather_idx_size_in; + m_gather_mode = gather_mode_in; + } + + + struct + { + //reserve the first 32 bits + union + { + struct + { + uint32_t m_count : 2; + uint32_t m_is_restore : 1; + uint32_t m_is_store : 1; + uint32_t m_nv : 1; + uint32_t m_scope_trait : 2; + uint32_t m_th :3; + uint32_t m_reserved_space : 20; + uint32_t m_gather_index_size : 1; + uint32_t m_gather_mode : 1; + }; + + uint32_t m_reserved0; + }; + + uint32_t m_lds_addr; + uint32_t m_global_addr_lo; + union + { + struct + { + uint32_t m_global_addr_hi : 25; + uint32_t m_reserved2 : 5; + uint32_t m_type : 2; + }; + uint32_t m_sgpr3; + }; + }; + __hip_uint32x4 m_bitfield; + + // setters for all user configurable fields + __device__ void count(bool value) + { + m_count = (int)value; + } + + __device__ void ldsAddr(uint32_t value) + { + m_lds_addr = value; + } + + __device__ void gatherMode(int enable, int value) + { + m_gather_mode = enable; + m_gather_index_size = value; + } + + __device__ void globalAddr(uintptr_t value) + { + m_global_addr_lo = value & 0xFFFFFFFF; + m_global_addr_hi = (value >> 32); + } +}; + +union gfx1250_TDM_GROUP1 +{ + __device__ constexpr gfx1250_TDM_GROUP1() : m_bitfield{0,0,0,0,0,0,0,0} {} + + struct + { + union + { + struct + { + uint32_t m_workgroup_mask : 16; + uint32_t m_data_size : 2; + uint32_t m_atomic_barrier_enable : 1; + uint32_t m_iterate_enable : 1; + uint32_t m_pad_enable : 1; + uint32_t m_early_timeout : 1; + uint32_t m_pad_interval : 3; + uint32_t m_pad_amount : 7; + }; + uint32_t m_sgpr0; + }; + union + { + struct + { + uint32_t m_atomic_barrier_address : 16; + uint32_t m_tensor_dim0_lo : 16; + }; + uint32_t m_sgpr1; + }; + union + { + struct + { + uint32_t m_tensor_dim0_hi : 16; + uint32_t m_tensor_dim1_lo : 16; + }; + uint32_t m_sgpr2; + }; + union + { + struct + { + uint32_t m_tensor_dim1_hi : 16; + uint32_t m_tile_dim0 : 16; + }; + uint32_t m_sgpr3; + }; + union + { + struct + { + uint32_t m_tile_dim1 : 16; + uint32_t m_tile_dim2 : 16; + }; + uint32_t m_sgpr4; + }; + union + { + uint32_t m_tensor_dim0_stride_lo; + uint32_t m_sgpr5; + }; + union + { + struct + { + uint32_t m_tensor_dim0_stride_hi : 16; + uint32_t m_tensor_dim1_stride_lo : 16; + }; + uint32_t m_sgpr6; + }; + union + { + uint32_t m_tensor_dim1_stride_hi; + uint32_t m_sgpr7; + }; + }; + __hip_uint32x8 m_bitfield; + + // setters for all fields + void __device__ inline workgroupMask(uint32_t value) + { + m_workgroup_mask = value; + } + + void __device__ inline dataSize(uint32_t value) + { + m_data_size = value; + } + + void __device__ inline atomicBarrierEnable(bool value) + { + m_atomic_barrier_enable = value; + } + + void __device__ inline iterateEnable(bool value) + { + m_iterate_enable = value; + } + + void __device__ inline padEnable(bool value) + { + m_pad_enable = value; + } + + void __device__ inline earlyTimeout(bool value) + { + m_early_timeout = value; + } + + void __device__ inline padInterval(uint32_t value) + { + m_pad_interval = value; + } + + void __device__ inline padAmount(uint32_t value) + { + m_pad_amount = value; + } + + void __device__ inline atomicBarrierAddress(uint32_t value) + { + m_atomic_barrier_address = value; + } + + void __device__ inline tileDim0(uint32_t value) + { + m_tile_dim0 = value; + } + + void __device__ inline tileDim1(uint32_t value) + { + m_tile_dim1 = value; + } + void __device__ inline tileDim2(uint32_t value) + { + m_tile_dim2 = value; + } + + void __device__ inline tensorDim0(uint32_t value) + { + m_tensor_dim0_lo = value & 0xFFFF; + m_tensor_dim0_hi = (value >> 16); + } + void __device__ inline tensorDim1(uint32_t value) + { + m_tensor_dim1_lo = value & 0xFFFF; + m_tensor_dim1_hi = (value >> 16); + } + void __device__ inline tensorDim0Stride(uint64_t value) + { + m_tensor_dim0_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim0_stride_hi = (value >> 32); + } + + void __device__ inline tensorDim1Stride(uint64_t value) + { + m_tensor_dim1_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim1_stride_hi = (value >> 32); + } +}; + +union gfx1250_TDM_GROUP2 +{ + __device__ gfx1250_TDM_GROUP2() : m_bitfield{0,0,0,0} {} + + struct + { + uint32_t m_tensor_dim2; //sgpr0 + uint32_t m_tensor_dim3; //sgpr1 + uint32_t m_tensor_dim2_stride_lo; //sgpr2 + union + { + struct + { + uint32_t m_tensor_dim2_stride_hi : 16; + uint32_t m_tile_dim3 : 16; + }; + uint32_t m_sgpr3; + }; + }; + __hip_uint32x4 m_bitfield; + + void __device__ inline tensorDim2Stride(uint64_t value) + { + m_tensor_dim2_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim2_stride_hi = value >> 32; + } +}; + +union gfx1250_TDM_GROUP3 +{ + __device__ gfx1250_TDM_GROUP3() : m_bitfield{0,0,0,0} {} + + struct + { + uint32_t m_tensor_dim3_stride_lo; //sgpr0 + union + { + struct + { + uint32_t m_tensor_dim3_stride_hi : 16; + uint32_t m_tensor_dim_4_lo : 16; + }; + uint32_t m_sgpr1; + }; + union + { + struct + { + uint32_t m_tensor_dim_4_hi : 16; + uint32_t m_tile_dim4 : 16; + }; + uint32_t m_sgpr2; + }; + uint32_t m_sgpr3_reserved; + }; + __hip_uint32x4 m_bitfield; + + void __device__ inline tensorDim3Stride(uint64_t value) + { + m_tensor_dim3_stride_lo = value & 0xFFFFFFFF; + m_tensor_dim3_stride_hi = value >> 32; + } + + void __device__ inline tensorDim4(uint32_t value) + { + m_tensor_dim_4_lo = value & 0xFFFF; + m_tensor_dim_4_hi = value >> 16; + } +}; +#pragma clang diagnostic pop +#endif // HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H From 9eede61c6a19bfe22e7b9be4db2185077d77e901 Mon Sep 17 00:00:00 2001 From: Tim <43156029+AtlantaPepsi@users.noreply.github.com> Date: Sun, 21 Jun 2026 23:04:52 -0500 Subject: [PATCH 2/5] Dynamic Loading IB verbs (#253) --- CMakeLists.txt | 87 +---- Makefile | 56 +-- build_packages_local.sh | 2 - src/client/Client.cpp | 4 +- src/client/EnvVars.hpp | 77 ++-- src/client/Topology.hpp | 3 +- src/header/TransferBench.hpp | 343 ++++++++--------- third-party/ibverbs/IbvDynLoad.hpp | 187 +++++++++ third-party/ibverbs/IbvHeader.hpp | 586 +++++++++++++++++++++++++++++ 9 files changed, 989 insertions(+), 356 deletions(-) create mode 100644 third-party/ibverbs/IbvDynLoad.hpp create mode 100644 third-party/ibverbs/IbvHeader.hpp 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 46858abc..1e52f026 100644 --- a/Makefile +++ b/Makefile @@ -19,9 +19,7 @@ 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) @@ -91,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 69bd23be..23377321 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -377,27 +377,23 @@ 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"); @@ -454,9 +450,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"); } @@ -510,32 +506,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")); @@ -543,10 +538,6 @@ 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"); 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/header/TransferBench.hpp b/src/header/TransferBench.hpp index d5072b1b..b81f5634 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -59,9 +59,7 @@ THE SOFTWARE. #include #include -#ifdef NIC_EXEC_ENABLED -#include -#endif +#include "IbvDynLoad.hpp" #ifdef MPI_COMM_ENABLED #include @@ -1686,16 +1684,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; @@ -1819,7 +1835,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}; @@ -2127,14 +2143,14 @@ 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"}); + 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"}); + } } -#endif // NVIDIA specific #if defined(__NVCC__) @@ -2527,7 +2543,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) { @@ -2581,11 +2597,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; } @@ -2822,7 +2837,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 @@ -2847,7 +2861,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 @@ -2905,7 +2918,6 @@ namespace { } }; -#ifdef NIC_EXEC_ENABLED // Structure to track information about IBV devices struct IbvDevice { @@ -2918,12 +2930,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]); @@ -3036,7 +3047,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; @@ -3121,9 +3132,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } return ibvDeviceList; } -#endif // NIC_EXEC_ENABLED -#ifdef NIC_EXEC_ENABLED // PCIe-related functions //======================================================================================== @@ -3308,9 +3317,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } return matches; } -#endif // NIC_EXEC_ENABLED -#ifdef NIC_EXEC_ENABLED // IB Verbs-related functions //======================================================================================== @@ -3481,16 +3488,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 } } @@ -3519,27 +3522,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()); @@ -3585,27 +3583,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()); @@ -3776,16 +3769,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) { @@ -3813,7 +3806,6 @@ static bool IsConfiguredGid(union ibv_gid const& gid) return ERR_NONE; } -#endif // NIC_EXEC_ENABLED // Data validation-related functions //======================================================================================== @@ -4504,14 +4496,14 @@ 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 @@ -4608,11 +4600,9 @@ 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 @@ -4740,7 +4730,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, @@ -4841,7 +4830,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid) } return ERR_NONE; } -#endif + // GFX Executor-related functions //======================================================================================== @@ -6100,9 +6089,7 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, case EXE_GPU_ASYNC_TENSOR: return RunGpuAsyncTensorExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); case EXE_GPU_ASYNC_MEMOPS: return RunGpuAsyncMemOpsExecutor(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 @@ -6509,11 +6496,7 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, 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; @@ -7762,22 +7745,23 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, // 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++) { @@ -7802,101 +7786,100 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, } // 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++) { @@ -7916,20 +7899,20 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, 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 } } 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" From d7be15647e25e85f7e0345b4d07f82e98fe2434f Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Wed, 8 Jul 2026 14:48:37 -0500 Subject: [PATCH 3/5] TDM productionization in TB --- src/client/EnvVars.hpp | 34 +- src/client/Presets/Help.hpp | 13 +- src/client/Utilities.hpp | 4 +- src/header/TransferBench.hpp | 637 ++++++++++++++--------------------- src/header/tdm.h | 355 ------------------- src/header/tdmCopy.h | 495 +++++++++++++++++++++++++++ 6 files changed, 785 insertions(+), 753 deletions(-) delete mode 100644 src/header/tdm.h create mode 100644 src/header/tdmCopy.h diff --git a/src/client/EnvVars.hpp b/src/client/EnvVars.hpp index 23377321..f69fc9b5 100644 --- a/src/client/EnvVars.hpp +++ b/src/client/EnvVars.hpp @@ -127,7 +127,9 @@ class EnvVars // 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 tdmPipelined; // Use the pipelined (depth=2) TDM tensor kernel (0=non-pipelined, 1=pipelined) + 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 @@ -180,7 +182,9 @@ class EnvVars showPercentiles = GetEnvVarArray("SHOW_PERCENTILES", {}); tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); tdmMaxLdsBytes = GetEnvVar("TDM_MAX_LDS_BYTES" , INT_MAX); - tdmPipelined = GetEnvVar("TDM_PIPELINED" , 0); + 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); @@ -398,10 +402,14 @@ class EnvVars 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(" 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_PIPELINED - 1=use pipelined (depth=2) TDM kernel, 0=use non-pipelined kernel\n"); - printf(" USE_HIP_EVENTS - Use HIP events for GFX executor timing\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"); @@ -543,16 +551,23 @@ class EnvVars "%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_PIPELINED", tdmPipelined, - "Using %s TDM tensor kernel", tdmPipelined ? "pipelined (depth=2)" : "non-pipelined"); + 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, @@ -761,7 +776,10 @@ class EnvVars cfg.tdm.blockSize = tdmBlockSize; cfg.tdm.maxLDSBytes = tdmMaxLdsBytes; - cfg.tdm.pipelined = tdmPipelined; + 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 de92ee9e..8049630c 100644 --- a/src/client/Presets/Help.hpp +++ b/src/client/Presets/Help.hpp @@ -37,15 +37,16 @@ int HelpPreset([[maybe_unused]] EnvVars& ev, printf("# SRC 1 -> Executor -> DST 1\n"); printf("# SRC X DST Y\n"); printf("\n"); - printf("# Seven Executors are supported by TransferBench\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) Batched-DMA Batch item (Must have single SRC, at least one DST)\n"); - printf("# 6) TDM GPU threadblock/Compute Unit (CU)\n"); - printf("# 7) Async Load/Store GPU threadblock/Compute Unit (CU)\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"); @@ -73,8 +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 kernel (Indexed from 0 to # GPUs - 1)\n"); - printf("# - L: GPU async load/store (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"); diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index 77a44b3c..af9a1696 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -439,8 +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_ASYNC_TENSOR: return "AT"; // async tensor kernel path - case EXE_GPU_ASYNC_MEMOPS: return "AL"; // async load/store kernel path + 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 b81f5634..7900f20d 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -25,7 +25,6 @@ THE SOFTWARE. #include #include #include -#include #include #include #include @@ -79,7 +78,7 @@ THE SOFTWARE. #ifdef AMD_SMI_ENABLED #include "amd_smi/amdsmi.h" #endif -#include "tdm.h" +#include "tdmCopy.h" #endif /// @endcond // Batched DMA executor is only supported with HIP >= 7.1 and CUDA 12.8 @@ -109,14 +108,14 @@ 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_ASYNC_TENSOR = 6, ///< GPU async kernel — tensor-op path (subExecutor — reserved; stub uses 1-D launch) - EXE_GPU_ASYNC_MEMOPS = 7, ///< GPU async kernel — load/store path (subExecutor — reserved; stub uses 1-D launch) + 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[9] = "CGDINBTL"; + 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 - || e == EXE_GPU_ASYNC_TENSOR || e == EXE_GPU_ASYNC_MEMOPS; + || e == EXE_GPU_TDM || e == EXE_GPU_ALS; } inline bool IsNicExeType(ExeType e){ return e == EXE_NIC || e == EXE_NIC_NEAREST; } @@ -284,9 +283,13 @@ namespace TransferBench struct TdmOptions { - int blockSize = 256; ///< Size of each threadblock (must be multiple of 64) - int maxLDSBytes = INT_MAX; ///< Maximum number of bytes of __shared__ memory to use - int pipelined = 0; ///< Whether to call the pipelined TDM kernel + int blockSize = 256; ///< Size of each threadblock (must be multiple of 64) + 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 }; @@ -2142,6 +2145,41 @@ namespace { } } + // 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)) { @@ -2275,8 +2313,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 || - t.exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || t.exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { + 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, @@ -2334,24 +2371,35 @@ namespace { hasFatalError = true; } break; - case EXE_GPU_ASYNC_TENSOR: - case EXE_GPU_ASYNC_MEMOPS: + 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: Async GPU kernel executor requires exactly 1 SRC and 1 DST", i}); + "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: Async GPU kernel device index must be between 0 and %d (instead of %d) for rank %d", + "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: Async GPU kernel executor does not support subindices yet", i}); + "Transfer %d: GPU TDM executor does not support subindices yet", i}); + hasFatalError = true; + break; + } + 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; } break; @@ -2375,7 +2423,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 } @@ -2725,18 +2772,18 @@ namespace { } break; } - case EXE_GPU_ASYNC_TENSOR: - case EXE_GPU_ASYNC_MEMOPS: + case EXE_GPU_TDM: +// case EXE_GPU_ALS: { int numGpuSubExec = GetNumSubExecutors(exeDevice); if (totalSubExecs[exeDevice] > numGpuSubExec) errors.push_back({ERR_WARN, - "GPU %d (async kernel) requests %d total subexecutors however only %d CUs available. " + "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 (async kernel) attempting %d parallel transfers, however GPU_MAX_HW_QUEUES only set to %d", + "GPU %d TDM attempting %d parallel transfers, however GPU_MAX_HW_QUEUES only set to %d", exeDevice.exeIndex, transferCount[exeDevice], gpuMaxHwQueues}); } break; @@ -4374,16 +4421,14 @@ namespace { } // Prepare additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || - exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) - && 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_ASYNC_TENSOR || - exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) + exeDevice.exeType == EXE_GPU_TDM || + exeDevice.exeType == EXE_GPU_ALS) ? exeInfo.resources.size() : 1; exeInfo.streams.resize(numStreamsToUse); @@ -4401,8 +4446,7 @@ namespace { } } - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents || exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || - exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { + 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) { @@ -4412,8 +4456,9 @@ namespace { } } - // 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__) @@ -4434,7 +4479,8 @@ namespace { 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; @@ -4506,8 +4552,9 @@ namespace { } } - // 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, @@ -4606,21 +4653,17 @@ namespace { } // Teardown additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || - exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) - && 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 || exeDevice.exeType == EXE_GPU_ASYNC_TENSOR || - exeDevice.exeType == EXE_GPU_ASYNC_MEMOPS) { - 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 @@ -4952,203 +4995,6 @@ namespace { } } - //---------------------------------------------------------------------------- -#if defined(__NVCC__) - __global__ void GpuAsyncTensorOpsStubKernel(float const* __restrict__ src, - float* __restrict__ dst, - size_t numFloats) - { - size_t const gid = blockIdx.x * blockDim.x + threadIdx.x; - size_t const stride = gridDim.x * blockDim.x; - for (size_t i = gid; i < numFloats; i += stride) - dst[i] = src[i]; - } -#else -// The TDM API does not have a function to set the transfer size, so we need to do it manually. -__device__ void SetTransferSize(gfx1250_TDM_GROUP1& group1, int numElements){ - group1.tensorDim0(numElements); - group1.tensorDim0Stride(numElements); - group1.tileDim0(numElements); -} -namespace { - __constant__ constexpr bool verbose = false; // Turn off to disable debug prints from kernel TDM kernels. -} - -// numElementsPerTile is the number of elements to process per tile. Its maximum value is the shared memory size / number of waves per workgroup. -__global__ void GpuAsyncTensorOpsKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerTile){ - extern __shared__ __align__(128) float shmem[]; - int waveId = threadIdx.x / warpSize; - int numWavesPerBlock = blockDim.x / warpSize; - size_t itemsProcessedPerGridIteration = numElementsPerTile * numWavesPerBlock * gridDim.x; - - float* shmemPtr = static_cast(shmem) + numElementsPerTile * waveId; - // Local per-wave source and destination pointers - const float* srcPtr = src + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); - float* dstPtr = dst + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); - gfx1250_TDM_GROUP0 group0; - - group0.ldsAddr((uintptr_t)shmemPtr); - - gfx1250_TDM_GROUP1 group1; - group1.dataSize(2); // Log2 of the element size in bytes, so 2 for float - SetTransferSize(group1, numElementsPerTile); - - constexpr __hip_uint32x4 empty_x4{}; - constexpr __hip_uint32x8 empty_x8{}; - size_t elementsToProcess = numElementsPerTile; - while(srcPtr < src + numElements){ - // Handle the last tile of the block, which may be less than num_elements_per_tile. - if(src + numElements - srcPtr < numElementsPerTile){ - size_t remainingElements = src + numElements - srcPtr; - SetTransferSize(group1, remainingElements); - if constexpr(verbose) elementsToProcess = remainingElements; - } - // Copy from global memory to LDS - group0.globalAddr((uintptr_t)srcPtr); - if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Block %d, Wave %d: Loading %zu elements from %p to shared memory at %p\n", blockIdx.x, waveId, elementsToProcess, srcPtr, shmemPtr); - __builtin_amdgcn_tensor_load_to_lds(group0.m_bitfield, group1.m_bitfield, empty_x4, empty_x4, empty_x8, 0); - __builtin_amdgcn_s_wait_tensorcnt(0); - - // write back from LDS to global - group0.globalAddr((uintptr_t)dstPtr); - if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Block %d, Wave %d: Storing %zu elements from shared memory at %p to %p\n", blockIdx.x, waveId, elementsToProcess, shmemPtr, dstPtr); - __builtin_amdgcn_tensor_store_from_lds(group0.m_bitfield, group1.m_bitfield, empty_x4, empty_x4, empty_x8, 0); - __builtin_amdgcn_s_wait_tensorcnt(0); - srcPtr += itemsProcessedPerGridIteration; - dstPtr += itemsProcessedPerGridIteration; - } -} - -struct TileMover { - gfx1250_TDM_GROUP0 group0; - gfx1250_TDM_GROUP1 group1; - float* dstPtr{nullptr}; - float* shmemPtr{nullptr}; // TODO: Remove this - int numElementsToProcess{0}; // TODO: Remove this - - __device__ void LoadTile(float* shmemPtr, const float* srcPtr, float* dstPtr, int numElementsToProcess) { - group0.ldsAddr((uintptr_t)shmemPtr); - group0.globalAddr((uintptr_t)srcPtr); - group1.dataSize(2); - SetTransferSize(group1, numElementsToProcess); - this -> dstPtr = dstPtr; this -> shmemPtr = shmemPtr; this -> numElementsToProcess = numElementsToProcess; - if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Loading %d elements from %p to %p\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess, srcPtr, shmemPtr); - __builtin_amdgcn_tensor_load_to_lds(group0.m_bitfield, group1.m_bitfield, __hip_uint32x4{}, __hip_uint32x4{}, __hip_uint32x8{}, 0); - } - __device__ void StoreTile() { - assert(dstPtr != nullptr); - group0.globalAddr((uintptr_t)dstPtr); - if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Storing %d elements from %p to %p\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess, shmemPtr, dstPtr); - __builtin_amdgcn_tensor_store_from_lds(group0.m_bitfield, group1.m_bitfield, __hip_uint32x4{}, __hip_uint32x4{}, __hip_uint32x8{}, 0); - } -}; - -// numElementsPerTile is the number of elements to process per tile. Its maximum value is the shared memory size / number of waves per workgroup. -__global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerSubtile){ - extern __shared__ __align__(128) float shmem[]; - int waveId = threadIdx.x / warpSize; - int numWavesPerBlock = blockDim.x / warpSize; // This only works because the blockDim.x is fixed. - constexpr int pipelineDepth = 2; - - int numElementsPerTile = pipelineDepth * numElementsPerSubtile; - size_t itemsProcessedPerGridIteration = numElementsPerTile * numWavesPerBlock * gridDim.x; - - float* shmemPtr = static_cast(shmem) + numElementsPerTile * waveId; - // Local per-wave source and destination pointers - const float* srcPtr = src + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); - float* dstPtr = dst + numElementsPerTile * (waveId + blockIdx.x * numWavesPerBlock); - - // Skip waves whose initial position is already past the end of the buffer. - if (srcPtr >= src + numElements){ - if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Warp %d Block %d: Skipping wave %d\n", threadIdx.x / warpSize, blockIdx.x, waveId); - return; - } - - int numElementsToProcess = std::min(numElementsPerSubtile, (int)(src + numElements - srcPtr)); - TileMover tileMovers[pipelineDepth]{}; - unsigned int slot = 0; - tileMovers[slot].LoadTile(shmemPtr, srcPtr, dstPtr, numElementsToProcess); - while(srcPtr + (slot ^ 1) * numElementsPerSubtile < src + numElements){ - slot ^= 1; - // Handle the last tile of the block, which may be less than num_elements_per_tile. - if(src + numElements - (srcPtr + slot * numElementsPerSubtile) < numElementsPerSubtile){ - numElementsToProcess = src + numElements - (srcPtr + slot * numElementsPerSubtile); - if constexpr(verbose) if (threadIdx.x % warpSize == 0) printf("Small Subtile: Warp %d Block %d: Num Elements To Process: %d\n", threadIdx.x / warpSize, blockIdx.x, numElementsToProcess); - } - // Copy from global memory to LDS - tileMovers[slot].LoadTile(shmemPtr + slot * numElementsPerSubtile, srcPtr + slot * numElementsPerSubtile, dstPtr + slot * numElementsPerSubtile, numElementsToProcess); - __builtin_amdgcn_s_wait_tensorcnt(1); - - // write back from LDS to global for the other slot - tileMovers[slot^1].StoreTile(); - __builtin_amdgcn_s_wait_tensorcnt(1); - srcPtr += itemsProcessedPerGridIteration * slot; - dstPtr += itemsProcessedPerGridIteration * slot; - } - __builtin_amdgcn_s_wait_tensorcnt(0); - tileMovers[slot].StoreTile(); - __builtin_amdgcn_s_wait_tensorcnt(0); -} -#endif - - // Source for definitions: https://github.com/llvm/llvm-project/blob/4e3bac3ea2cc6fd778d53e317dd9fc27c1ddfc4f/clang/include/clang/Basic/BuiltinsAMDGPU.td#L956 - // and internal ISA documentation - using uint32x4 = __attribute__((__vector_size__(4 * sizeof(int)))) int; - // Loads 16 bytes/lane from global (src) into LDS (dst). - __device__ void asyncLoadX4(float const* src, float* dst){ - __builtin_amdgcn_global_load_async_to_lds_b128( - (__attribute__((address_space(1))) uint32x4*)src, - (__attribute__((address_space(3))) uint32x4*)dst, 0, 0); - } - // Stores 16 bytes/lane from LDS (src) to global (dst). - __device__ void asyncStoreX4(float const* src, float* dst){ - __builtin_amdgcn_global_store_async_from_lds_b128( - (__attribute__((address_space(1))) uint32x4*)dst, - (__attribute__((address_space(3))) uint32x4*)src, 0, 0); - } - - // numElementsPerTile should be a multiple of 128, given the asyncLoadX4 and asyncStoreX4 functions operate on a 32-lane warp and 4 elements per lane (512 bytes total) - __global__ void GpuAsyncLoadStoreKernel(float const* __restrict__ src, float* __restrict__ dst, size_t numElements, int numElementsPerTile) - { - extern __shared__ __align__(128) float shmem[]; - int waveId = threadIdx.x / warpSize; - int laneId = threadIdx.x % warpSize; - int numWavesPerBlock = blockDim.x / warpSize; - size_t itemsProcessedPerGridIteration = (size_t)numElementsPerTile * numWavesPerBlock * gridDim.x; - constexpr size_t stride_per_lane = sizeof(float4) / sizeof(float); // 4 elements per lane - // Number of elements a full 32-lane warp moves per b128 async op (32 * 4 = 128). - size_t elements_per_dwordx4_load_store = warpSize * stride_per_lane; - - float* shmemPtr = shmem + numElementsPerTile * waveId + laneId * stride_per_lane; - // Wave-tile base (lane-independent) used to gate the loop and compute the partial-tile size. - size_t tileBase = (size_t)numElementsPerTile * (waveId + (size_t)blockIdx.x * numWavesPerBlock); - - while(tileBase < numElements){ - // Handle the last tile, which may be less than numElementsPerTile. - size_t elementsToProcess = numElementsPerTile; - if(numElements - tileBase < (size_t)numElementsPerTile) - elementsToProcess = numElements - tileBase; - - const float* srcPtr = src + tileBase + laneId * stride_per_lane; - float* dstPtr = dst + tileBase + laneId * stride_per_lane; - - // Copy from global memory to LDS - for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store) { - asyncLoadX4(srcPtr + i, shmemPtr + i); - __builtin_amdgcn_s_wait_asynccnt(0); - } - - // Write back from LDS to global - for (size_t i = 0; i < elementsToProcess; i += elements_per_dwordx4_load_store){ - asyncStoreX4(shmemPtr + i, dstPtr + i); - __builtin_amdgcn_s_wait_asynccnt(0); - } - - tileBase += itemsProcessedPerGridIteration; - } -} - - // Simplified Kernel for GFX execution for copies only template __global__ void __launch_bounds__(LAUNCH_BOUND) @@ -5736,132 +5582,6 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, return ERR_NONE; } - static ErrResult ExecuteGpuAsyncKernelTransfer(int const iteration, - ConfigOptions const& cfg, - int const exeIndex, - hipStream_t const stream, - hipEvent_t const startEvent, - hipEvent_t const stopEvent, - bool const tensorFlavor, - int const numSubExecs, - TransferResources& rss) - { - ERR_CHECK(hipSetDevice(exeIndex)); - int const initOffset = cfg.data.byteOffset / sizeof(float); - size_t const numFloats = rss.numBytes / sizeof(float); - float const* __restrict__ src = rss.srcMem[0] + initOffset; - float* __restrict__ dst = rss.dstMem[0] + initOffset; - - int const threads = tensorFlavor ? cfg.tdm.blockSize : cfg.gfx.blockSize; - int const blocks = numSubExecs; - - auto cpuStart = std::chrono::high_resolution_clock::now(); - - int subIterations = 0; - do { - if (startEvent) - ERR_CHECK(hipEventRecord(startEvent, stream)); -#if defined(__NVCC__) - GpuAsyncTensorOpsStubKernel<<>>(src, dst, numFloats); -#else - bool usePipelinedTensorOps = (cfg.tdm.pipelined != 0); - auto gpuKernel = tensorFlavor ? (usePipelinedTensorOps ? GpuAsyncPipelinedTensorOpsKernel : GpuAsyncTensorOpsKernel) : GpuAsyncLoadStoreKernel; - if (rss.numBytes % sizeof(float) != 0) - return {ERR_FATAL, "Async tensor executor (TDM): numBytes (%zu) must be a multiple of %zu", rss.numBytes, sizeof(float)}; - - - hipDeviceProp_t props{}; - ERR_CHECK(hipGetDeviceProperties(&props, exeIndex)); - int const warpSize = props.warpSize; - int maxShmem = 0; - ERR_CHECK(hipDeviceGetAttribute(&maxShmem, hipDeviceAttributeMaxSharedMemoryPerBlock, exeIndex)); - maxShmem = std::min(maxShmem, cfg.tdm.maxLDSBytes); - int kWavesPerWorkgroup = threads / warpSize; - int pipelineDepth = usePipelinedTensorOps ? 2 : 1; - int numFloatsPerTile = std::max(1, maxShmem / (sizeof(float) * kWavesPerWorkgroup * pipelineDepth)); - unsigned int shmemBytes = numFloatsPerTile * kWavesPerWorkgroup * sizeof(float) * pipelineDepth; - hipLaunchKernelGGL(gpuKernel, dim3(blocks), dim3(threads), - shmemBytes, stream, src, dst, numFloats, - numFloatsPerTile); -#endif - ERR_CHECK(hipGetLastError()); - if (stopEvent) - ERR_CHECK(hipEventRecord(stopEvent, stream)); - ERR_CHECK(hipStreamSynchronize(stream)); - } while (++subIterations != cfg.general.numSubIterations); - - if (iteration >= 0) { - double deltaMsec; - if (startEvent && stopEvent) { - float gpuDeltaMsec; - ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); - deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations; - } else { - auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; - deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 - / cfg.general.numSubIterations; - } - rss.totalDurationMsec += deltaMsec; - if (cfg.general.recordPerIteration) - rss.perIterMsec.push_back(deltaMsec); - } - return ERR_NONE; - } - - static ErrResult RunGpuAsyncKernelExecutor(int const iteration, - ConfigOptions const& cfg, - int const exeIndex, - bool const tensorFlavor, - 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 = (!exeInfo.startEvents.empty()) ? exeInfo.startEvents[i] : nullptr; - hipEvent_t stopEv = (!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, - ExecuteGpuAsyncKernelTransfer, - iteration, - std::cref(cfg), - exeIndex, - exeInfo.streams[i], - startEv, - stopEv, - tensorFlavor, - 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; - } - - static ErrResult RunGpuAsyncTensorExecutor(int const iteration, - ConfigOptions const& cfg, - int const exeIndex, - ExeInfo& exeInfo) - { - return RunGpuAsyncKernelExecutor(iteration, cfg, exeIndex, true, exeInfo); - } - - static ErrResult RunGpuAsyncMemOpsExecutor(int const iteration, - ConfigOptions const& cfg, - int const exeIndex, - ExeInfo& exeInfo) - { - return RunGpuAsyncKernelExecutor(iteration, cfg, exeIndex, false, exeInfo); - } - // DMA Executor-related functions //======================================================================================== @@ -6076,6 +5796,160 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, } #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); + 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, @@ -6086,8 +5960,7 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, 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_ASYNC_TENSOR: return RunGpuAsyncTensorExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); - case EXE_GPU_ASYNC_MEMOPS: return RunGpuAsyncMemOpsExecutor(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); case EXE_NIC: return RunNicExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); #ifdef BMA_EXEC_ENABLED @@ -6782,8 +6655,8 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, result |= RecursiveWildcardTransferExpansion(wc, baseRankIndex, numBytes, numSubExecs, transfers); wc.exe.exeSubIndices[0] = -2; return result; - case EXE_GPU_ASYNC_TENSOR: - case EXE_GPU_ASYNC_MEMOPS: + 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; @@ -7672,11 +7545,11 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, 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_ASYNC_TENSOR] = numGpus; - topo.numExecutors[EXE_GPU_ASYNC_MEMOPS] = numGpus; + topo.numExecutors[EXE_GPU_TDM] = numGpus; + topo.numExecutors[EXE_GPU_ALS] = numGpus; std::vector gpuArchNames(numGpus); @@ -7700,8 +7573,8 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, 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_ASYNC_TENSOR, exeIndex}] = gpuName + " [async tensor]"; - topo.executorName[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = gpuName + " [async LD/ST]"; + 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]; @@ -7731,13 +7604,13 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, 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_ASYNC_TENSOR, exeIndex}] = 0; - topo.numExecutorSubIndices[{EXE_GPU_ASYNC_MEMOPS, 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_ASYNC_TENSOR, exeIndex}] = numDeviceCUs; - topo.numSubExecutors[{EXE_GPU_ASYNC_MEMOPS, exeIndex}] = numDeviceCUs; + topo.numSubExecutors[{EXE_GPU_TDM, exeIndex}] = numDeviceCUs; + topo.numSubExecutors[{EXE_GPU_ALS, exeIndex}] = numDeviceCUs; topo.closestCpuNumaToGpu[exeIndex] = closestNuma; topo.closestNicsToGpu[exeIndex] = {}; @@ -8158,7 +8031,7 @@ __global__ void GpuAsyncPipelinedTensorOpsKernel(float const* __restrict__ src, agent = cpuAgents[exeDevice.exeIndex]; break; case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: - case EXE_GPU_ASYNC_TENSOR: case EXE_GPU_ASYNC_MEMOPS: + 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]; diff --git a/src/header/tdm.h b/src/header/tdm.h deleted file mode 100644 index c55a2bdd..00000000 --- a/src/header/tdm.h +++ /dev/null @@ -1,355 +0,0 @@ -/* -Copyright (c) 2025 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. -*/ - -#ifndef HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H -#define HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H - -#if !defined(__HIPCC_RTC__) -/* hip_runtime_api.h defines hipMemoryType before including driver_types.h; do not include - * driver_types.h alone or HIP 7.x fails with unknown type hipMemoryType. */ -#include -#include -#if __cplusplus -#include -#else -#include -#endif -#endif - -using __hip_uint32x4 = __NATIVE_VECTOR__(4, uint32_t); -using __hip_uint32x8 = __NATIVE_VECTOR__(8, uint32_t); - -#pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wshift-count-overflow" // we're going to be using overflow on purpose - -union gfx1250_TDM_GROUP0 -{ - __device__ constexpr gfx1250_TDM_GROUP0() : m_bitfield{0x1, 0, 0, 0x80000000} - {} - - __device__ gfx1250_TDM_GROUP0( uintptr_t lds_addr_in, - uintptr_t global_addr_in, - uint32_t gather_idx_size_in = 0, - uint32_t gather_mode_in = 0 ) - { - globalAddr(global_addr_in); - m_count = 1; - m_lds_addr = lds_addr_in; - m_type = 2; - m_gather_index_size = gather_idx_size_in; - m_gather_mode = gather_mode_in; - } - - - struct - { - //reserve the first 32 bits - union - { - struct - { - uint32_t m_count : 2; - uint32_t m_is_restore : 1; - uint32_t m_is_store : 1; - uint32_t m_nv : 1; - uint32_t m_scope_trait : 2; - uint32_t m_th :3; - uint32_t m_reserved_space : 20; - uint32_t m_gather_index_size : 1; - uint32_t m_gather_mode : 1; - }; - - uint32_t m_reserved0; - }; - - uint32_t m_lds_addr; - uint32_t m_global_addr_lo; - union - { - struct - { - uint32_t m_global_addr_hi : 25; - uint32_t m_reserved2 : 5; - uint32_t m_type : 2; - }; - uint32_t m_sgpr3; - }; - }; - __hip_uint32x4 m_bitfield; - - // setters for all user configurable fields - __device__ void count(bool value) - { - m_count = (int)value; - } - - __device__ void ldsAddr(uint32_t value) - { - m_lds_addr = value; - } - - __device__ void gatherMode(int enable, int value) - { - m_gather_mode = enable; - m_gather_index_size = value; - } - - __device__ void globalAddr(uintptr_t value) - { - m_global_addr_lo = value & 0xFFFFFFFF; - m_global_addr_hi = (value >> 32); - } -}; - -union gfx1250_TDM_GROUP1 -{ - __device__ constexpr gfx1250_TDM_GROUP1() : m_bitfield{0,0,0,0,0,0,0,0} {} - - struct - { - union - { - struct - { - uint32_t m_workgroup_mask : 16; - uint32_t m_data_size : 2; - uint32_t m_atomic_barrier_enable : 1; - uint32_t m_iterate_enable : 1; - uint32_t m_pad_enable : 1; - uint32_t m_early_timeout : 1; - uint32_t m_pad_interval : 3; - uint32_t m_pad_amount : 7; - }; - uint32_t m_sgpr0; - }; - union - { - struct - { - uint32_t m_atomic_barrier_address : 16; - uint32_t m_tensor_dim0_lo : 16; - }; - uint32_t m_sgpr1; - }; - union - { - struct - { - uint32_t m_tensor_dim0_hi : 16; - uint32_t m_tensor_dim1_lo : 16; - }; - uint32_t m_sgpr2; - }; - union - { - struct - { - uint32_t m_tensor_dim1_hi : 16; - uint32_t m_tile_dim0 : 16; - }; - uint32_t m_sgpr3; - }; - union - { - struct - { - uint32_t m_tile_dim1 : 16; - uint32_t m_tile_dim2 : 16; - }; - uint32_t m_sgpr4; - }; - union - { - uint32_t m_tensor_dim0_stride_lo; - uint32_t m_sgpr5; - }; - union - { - struct - { - uint32_t m_tensor_dim0_stride_hi : 16; - uint32_t m_tensor_dim1_stride_lo : 16; - }; - uint32_t m_sgpr6; - }; - union - { - uint32_t m_tensor_dim1_stride_hi; - uint32_t m_sgpr7; - }; - }; - __hip_uint32x8 m_bitfield; - - // setters for all fields - void __device__ inline workgroupMask(uint32_t value) - { - m_workgroup_mask = value; - } - - void __device__ inline dataSize(uint32_t value) - { - m_data_size = value; - } - - void __device__ inline atomicBarrierEnable(bool value) - { - m_atomic_barrier_enable = value; - } - - void __device__ inline iterateEnable(bool value) - { - m_iterate_enable = value; - } - - void __device__ inline padEnable(bool value) - { - m_pad_enable = value; - } - - void __device__ inline earlyTimeout(bool value) - { - m_early_timeout = value; - } - - void __device__ inline padInterval(uint32_t value) - { - m_pad_interval = value; - } - - void __device__ inline padAmount(uint32_t value) - { - m_pad_amount = value; - } - - void __device__ inline atomicBarrierAddress(uint32_t value) - { - m_atomic_barrier_address = value; - } - - void __device__ inline tileDim0(uint32_t value) - { - m_tile_dim0 = value; - } - - void __device__ inline tileDim1(uint32_t value) - { - m_tile_dim1 = value; - } - void __device__ inline tileDim2(uint32_t value) - { - m_tile_dim2 = value; - } - - void __device__ inline tensorDim0(uint32_t value) - { - m_tensor_dim0_lo = value & 0xFFFF; - m_tensor_dim0_hi = (value >> 16); - } - void __device__ inline tensorDim1(uint32_t value) - { - m_tensor_dim1_lo = value & 0xFFFF; - m_tensor_dim1_hi = (value >> 16); - } - void __device__ inline tensorDim0Stride(uint64_t value) - { - m_tensor_dim0_stride_lo = value & 0xFFFFFFFF; - m_tensor_dim0_stride_hi = (value >> 32); - } - - void __device__ inline tensorDim1Stride(uint64_t value) - { - m_tensor_dim1_stride_lo = value & 0xFFFFFFFF; - m_tensor_dim1_stride_hi = (value >> 32); - } -}; - -union gfx1250_TDM_GROUP2 -{ - __device__ gfx1250_TDM_GROUP2() : m_bitfield{0,0,0,0} {} - - struct - { - uint32_t m_tensor_dim2; //sgpr0 - uint32_t m_tensor_dim3; //sgpr1 - uint32_t m_tensor_dim2_stride_lo; //sgpr2 - union - { - struct - { - uint32_t m_tensor_dim2_stride_hi : 16; - uint32_t m_tile_dim3 : 16; - }; - uint32_t m_sgpr3; - }; - }; - __hip_uint32x4 m_bitfield; - - void __device__ inline tensorDim2Stride(uint64_t value) - { - m_tensor_dim2_stride_lo = value & 0xFFFFFFFF; - m_tensor_dim2_stride_hi = value >> 32; - } -}; - -union gfx1250_TDM_GROUP3 -{ - __device__ gfx1250_TDM_GROUP3() : m_bitfield{0,0,0,0} {} - - struct - { - uint32_t m_tensor_dim3_stride_lo; //sgpr0 - union - { - struct - { - uint32_t m_tensor_dim3_stride_hi : 16; - uint32_t m_tensor_dim_4_lo : 16; - }; - uint32_t m_sgpr1; - }; - union - { - struct - { - uint32_t m_tensor_dim_4_hi : 16; - uint32_t m_tile_dim4 : 16; - }; - uint32_t m_sgpr2; - }; - uint32_t m_sgpr3_reserved; - }; - __hip_uint32x4 m_bitfield; - - void __device__ inline tensorDim3Stride(uint64_t value) - { - m_tensor_dim3_stride_lo = value & 0xFFFFFFFF; - m_tensor_dim3_stride_hi = value >> 32; - } - - void __device__ inline tensorDim4(uint32_t value) - { - m_tensor_dim_4_lo = value & 0xFFFF; - m_tensor_dim_4_hi = value >> 16; - } -}; -#pragma clang diagnostic pop -#endif // HIP_INCLUDE_HIP_AMD_GFX1250_TDM_H diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h new file mode 100644 index 00000000..48c49f66 --- /dev/null +++ b/src/header/tdmCopy.h @@ -0,0 +1,495 @@ +/* +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 +/// congtigous 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 +#if defined(__gfx1250__) && \ + __has_builtin(__builtin_amdgcn_tensor_load_to_lds) && \ + __has_include() + /* 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; } + return *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 From 33d0bf7a0d5dc199cb7d3aed13bff1454a1e5aba Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Mon, 13 Jul 2026 15:45:11 -0500 Subject: [PATCH 4/5] temporary fixes --- src/client/Presets/Help.hpp | 4 ++-- src/header/TransferBench.hpp | 13 ++++++++++++- src/header/tdmCopy.h | 2 +- 3 files changed, 15 insertions(+), 4 deletions(-) diff --git a/src/client/Presets/Help.hpp b/src/client/Presets/Help.hpp index 8049630c..e80987d9 100644 --- a/src/client/Presets/Help.hpp +++ b/src/client/Presets/Help.hpp @@ -113,11 +113,11 @@ 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 tensor-op path\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->L0->G1)\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"); diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 7900f20d..b69c1812 100644 --- a/src/header/TransferBench.hpp +++ b/src/header/TransferBench.hpp @@ -283,7 +283,7 @@ namespace TransferBench struct TdmOptions { - int blockSize = 256; ///< Size of each threadblock (must be multiple of 64) + 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) @@ -638,6 +638,7 @@ namespace TransferBench // Enumerations #define hipDeviceAttributeClockRate cudaDevAttrClockRate #define hipDeviceAttributeMultiprocessorCount cudaDevAttrMultiProcessorCount + #define hipDeviceAttributeMaxSharedMemoryPerBlock cudaDevAttrMaxSharedMemoryPerBlock #define hipDeviceAttributeWarpSize cudaDevAttrWarpSize #define hipErrorPeerAccessAlreadyEnabled cudaErrorPeerAccessAlreadyEnabled #define hipFuncCachePreferShared cudaFuncCachePreferShared @@ -667,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 @@ -2396,12 +2398,18 @@ namespace { 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) { @@ -5833,6 +5841,7 @@ namespace { int subIterations = 0; while (1) { tdm::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes); + __syncthreads(); if (++subIterations == numSubIterations) break; } @@ -8407,6 +8416,7 @@ namespace { // Enumerations #undef hipDeviceAttributeClockRate #undef hipDeviceAttributeMultiprocessorCount +#undef hipDeviceAttributeMaxSharedMemoryPerBlock #undef hipDeviceAttributeWarpSize #undef hipErrorPeerAccessAlreadyEnabled #undef hipFuncCachePreferShared @@ -8437,6 +8447,7 @@ namespace { #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 index 48c49f66..0c79875c 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -23,7 +23,7 @@ THE SOFTWARE. /// /// This is introduced via __device__ level memcpy-like API which can be /// either blocking or asynchronous, and utilize all warps, or a team of -/// congtigous warps, to allow for other warps to do other tasks +/// contiguous warps, to allow for other warps to do other tasks /// /// \par Quick start /// \code From d2aff82c5490bca9dee3009a9ad4a0b85306b4c8 Mon Sep 17 00:00:00 2001 From: AtlantaPepsi Date: Mon, 13 Jul 2026 21:21:14 -0500 Subject: [PATCH 5/5] gate runtime check for absent TDM header --- src/header/tdmCopy.h | 21 +++++++++++++++++++-- 1 file changed, 19 insertions(+), 2 deletions(-) diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h index 0c79875c..5175f8c6 100644 --- a/src/header/tdmCopy.h +++ b/src/header/tdmCopy.h @@ -77,9 +77,23 @@ THE SOFTWARE. #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) && \ - __has_include() + TDM_TOOLCHAIN_AVAILABLE /* extend: || (defined(__gfxNNNN__) && ...) */ # define TDM_SUPPORTED 1 #else @@ -133,7 +147,10 @@ __host__ __device__ inline bool IsTdmCopySupported(int deviceId = 0) { const char* arch = prop.gcnArchName; const char* p = "gfx1250"; while (*p && *arch == *p) { ++arch; ++p; } - return *p == '\0'; + // 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 }