diff --git a/CHANGELOG.md b/CHANGELOG.md index a102a52b..e1267b42 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,29 @@ Documentation for TransferBench is available at [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench). +## v1.70.00 +### Added +- Adding support for Tensor Data Mover (TDM)-based executor [T] on supported hardware (gfx1250, NVIDIA sm_90+ via TMA). + This provides an alternative data movement mechanism which utilizes async loads to shared memory / from shared memory +- Added support for SWEEP_MIN_POW2 and SWEEP_MAX_POW2 to set sweep bounds when bytes to transfer is 0 +- Added new "tdmsweep" preset that sweeps TDM executor options (block size / LDS / block order / subExecs) +- Added TB_SEND_USLEEP to insert a configurable microsecond delay after each socket SendData call (default: 0); useful for diagnosing small-message timing issues on sensitive clusters +- Adding ppodId / vpodId printing to verbose mode +- Added VALIDATE_ON_DEVICE to validate GPU destination (and source) memory via an on-device kernel instead of + copying back to the host. Expected values are pre-uploaded during preparation; only a mismatch count and the + first mismatch offset are returned. Takes precedence over VALIDATE_DIRECT for GPU destinations. +### Modified +- a2a, p2p, rings, poda2a, and podp2p presets now support the TDM executor +- CPU NUMA nodes with 0 cores will now be hidden. To re-enable, set TB_SHOW_ALL_NUMA=1 +- Switching to use of persistent threadpools to cut-down on thread creation overheads +- Updating default GFX unroll on GFX1250 to 32 +- Improved socket communicator robustness (TCP_NODELAY, partial send/recv handling, MSG_NOSIGNAL) +- Improved mismatch logging and smoketest fail reporting +- Destination memory is now cleared after each iteration when ALWAYS_VALIDATE is enabled, so each iteration starts from a known-zero state +### Fixed +- Guard before ibv_free_device_list to avoid invalid free +- Fix NIC to GPU proximity detection on systems with multiple PCIe domains + ## v1.69.01 ### Added - Added support for ABI change introduced in amd-smi 27.0.0 (ROCm 10.0) @@ -21,7 +44,7 @@ Documentation for TransferBench is available at - Created a top level third-party/ folder for ibverbs related files. Will also harbor future external source which TransferBench depends - Created a separate minimal header IbvHeader.hpp for ib verbs structs and IbvDynLoad.hpp for dynamic loading and status report for ib verbs functionality. - Dynamic loading is a singleton and done once per process, and TransferBench header will probe in runtime if basic ibverbs function as well as dmabuf export is supported. - - Also got rid HAVE_DMABUF_SUPPORT macro. Got rid of redundant dependency check on hsa header and rocr binaries (they are mandatory for AMD platform) in build process. Similar to ibv, it now dynamically checks for hsa_amd_portable_export_dmabuf symbol as part of check kernel support, and returns dmabuf support in runtime. + - Also got rid of HAVE_DMABUF_SUPPORT macro. Got rid of redundant dependency check on hsa header and rocr binaries (they are mandatory for AMD platform) in build process. Similar to ibv, it now dynamically checks for hsa_amd_portable_export_dmabuf symbol as part of check kernel support, and returns dmabuf support in runtime. ## v1.68.00 ### Fixed diff --git a/CMakeLists.txt b/CMakeLists.txt index fe585526..de1361fd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -103,8 +103,8 @@ set(ENV{ROCM_PATH} "${ROCM_PATH}") # TransferBench project definitions #================================================================================================== set(TRANSFERBENCH_VERSION_MAJOR 1) -set(TRANSFERBENCH_VERSION_MINOR 69) -set(TRANSFERBENCH_VERSION_PATCH_FALLBACK "01") +set(TRANSFERBENCH_VERSION_MINOR 70) +set(TRANSFERBENCH_VERSION_PATCH_FALLBACK "00") # Auto-compute patch from git: count commits since the last v..* tag. # Falls back to TRANSFERBENCH_VERSION_PATCH_FALLBACK when git is unavailable, diff --git a/Makefile b/Makefile index 37feae11..755598aa 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/docs/conceptual/transferbench-data-validation.rst b/docs/conceptual/transferbench-data-validation.rst index 815cfe90..4d60e749 100644 --- a/docs/conceptual/transferbench-data-validation.rst +++ b/docs/conceptual/transferbench-data-validation.rst @@ -142,6 +142,12 @@ environment variables or in a configuration file. - To compare GPU DST directly, set to ``1``. Supported on AMD hardware only and requires no host copy. To copy to host and compare, set to ``0``. + * - ``validateOnDevice`` + - ``VALIDATE_ON_DEVICE`` + - To compare GPU memory on the device via a kernel (expected values pre-uploaded during prep), set to ``1``. + Avoids the device-to-host copy; returns only a mismatch count and the first mismatch offset. + Takes precedence over ``validateDirect`` for GPU destinations. + * - ``validateSource`` - ``VALIDATE_SOURCE`` - To validate the SRC memory right after it's initialized, set to ``1`` (optional early check). diff --git a/docs/how to/use-transferbench.rst b/docs/how to/use-transferbench.rst index 21c8e00d..c34bc357 100644 --- a/docs/how to/use-transferbench.rst +++ b/docs/how to/use-transferbench.rst @@ -10,11 +10,10 @@ Using TransferBench You can control the SRC and DST memory locations by indicating the memory type followed by the device index. TransferBench supports the following memory types: -* Coarse-grained pinned host +* Pinned host (default, closest-GPU, coherent, non-coherent, and uncached) * Unpinned host -* Fine-grained host -* Coarse-grained global device -* Fine-grained global device +* Coarse-grained, fine-grained, and uncached global device +* Managed device * Null (for an empty transfer) In addition, you can determine the size of the transfer (number of bytes to copy) for the tests. @@ -120,12 +119,20 @@ Here is the list of arguments used to specify transfers in the config file: | Memory locations are specified by one or more device characters or device index pairs. | Characters indicate memory type and are followed by device index (0-indexed). | Here are the characters and their respective memory locations: - | - C: Pinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) - | - U: Unpinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) - | - B: Fine-grain host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) - | - G: Global device memory (on GPU device, indexed from 0 to [GPUs - 1]) - | - F: Fine-grain device memory (on GPU device, indexed from 0 to [GPUs - 1]) - | - N: Null memory (index ignored) + | - C: Pinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) + | - P: Pinned host memory (indexed by closest GPU, 0 to [GPUs-1]) + | - B: Coherent pinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) + | - D: Non-coherent pinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) + | - K: Uncached pinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) + | - H: Unpinned host memory (on NUMA node, indexed from 0 to [NUMA nodes-1]) + | - G: Global device memory (on GPU device, indexed from 0 to [GPUs-1]) + | - F: Fine-grain device memory (on GPU device, indexed from 0 to [GPUs-1]) + | - U: Uncached device memory (on GPU device, indexed from 0 to [GPUs-1]) + | - M: Managed device memory (on GPU device, indexed from 0 to [GPUs-1]) + | - N: Null memory (index ignored) + | + | ``D`` in an executor position is the DMA executor. ``D`` in a SRC/DST position is + | non-coherent pinned host memory. Round brackets and arrows "->" can be included for human clarity, but will be ignored. Lines starting with # are ignored while lines starting with ## are echoed to the output. @@ -140,6 +147,10 @@ Single DMA-executed transfer between GPU 0 and 1:: 1 1 (G0->D0->G1) +Unpinned host to GPU 0 using the DMA executor (PCIe H2D):: + + 1 1 (H0->D0->G0) + Copying 1Mb from GPU 0 to GPU 1 with 4 CUs, and 2Mb from GPU 1 to GPU 0 with 8 CUs:: -2 (G0->G0->G1 4 1M) (G1->G1->G0 8 2M) diff --git a/docs/reference/environment-variables.rst b/docs/reference/environment-variables.rst index 1b87f4df..7228e6e9 100644 --- a/docs/reference/environment-variables.rst +++ b/docs/reference/environment-variables.rst @@ -139,6 +139,12 @@ Data and validation options On AMD hardware, the CPU can directly access GPU device memory, avoiding the need for a staging buffer. This feature is not supported on NVIDIA hardware. - ``0`` + * - ``VALIDATE_ON_DEVICE`` + - Specifies whether to validate GPU memory on the device instead of copying it back to the host. Set to ``1`` to validate on-device, ``0`` to copy to host and compare. + + The expected values are pre-uploaded to the GPU during preparation, and a comparison kernel checks the destination in place, avoiding the device-to-host copy. Only a small result (mismatch count and first mismatch offset) is returned. Takes precedence over ``VALIDATE_DIRECT`` for GPU destinations; CPU destinations are always compared on the host. + - ``0`` + * - ``VALIDATE_SOURCE`` - Specifies whether to validate the source immediately after preparation. Set to ``1`` to validate, ``0`` to skip. diff --git a/src/client/Client.cpp b/src/client/Client.cpp index 5df1d62b..93e39954 100644 --- a/src/client/Client.cpp +++ b/src/client/Client.cpp @@ -171,7 +171,7 @@ int main(int argc, char **argv) } // Run the specified numbers of bytes otherwise generate a range of values - for (size_t bytes = (1<<10); bytes <= (1<<29); bytes *= 2) { + for (size_t bytes = (1ULL< fillPattern; // Pattern of floats used to fill source data vector fillCompress; // Percentages of 64B lines to be filled by random/1B0/2B0/4B0/32B0 + int sweepMaxPow2; // Maximum power of two to sweep up to when number of bytes to transfer is set to 0 + int sweepMinPow2; // Minimum power of two to sweep up from when number of bytes to transfer is set to 0 int validateDirect; // Validate GPU destination memory directly instead of staging GPU memory on host + int validateOnDevice; // Validate GPU dst/src memory via an on-device kernel instead of host memcmp int validateSource; // Validate source GPU memory immediately after preparation // DMA options @@ -124,6 +127,11 @@ class EnvVars int nicTrafficClass; // DSCP/traffic class byte for RoCE GRH int roceVersion; // RoCE version number + // TDM options + int tdmBlockOrder; // How threadblocks for multiple Transfers are ordered 0=sequential 1=interleaved + int tdmBlockSize; // Size of each threadblock for TDM kernels (must be multiple of 32) + int tdmLdsBytes; // Size of LDS (shared memory) bytes per threadblock for TDM kernels (0 = use device max) + // Developer features int gpuMaxHwQueues; // Tracks GPU_MAX_HW_QUEUES environment variable @@ -144,10 +152,11 @@ class EnvVars // Different hardware pick different GPU kernels // This performance difference is generally only noticable when executing fewer CUs int defaultGfxUnroll = 4; - if (archName == "gfx906") defaultGfxUnroll = 8; - else if (archName == "gfx90a") defaultGfxUnroll = 8; - else if (archName == "gfx942") defaultGfxUnroll = 4; - else if (archName == "gfx950") defaultGfxUnroll = 4; + if (archName == "gfx906") defaultGfxUnroll = 8; + else if (archName == "gfx90a") defaultGfxUnroll = 8; + else if (archName == "gfx942") defaultGfxUnroll = 4; + else if (archName == "gfx950") defaultGfxUnroll = 4; + else if (archName == "gfx1250") defaultGfxUnroll = 32; alwaysValidate = GetEnvVar("ALWAYS_VALIDATE", 0); blockBytes = GetEnvVar("BLOCK_BYTES" , 256); @@ -173,11 +182,20 @@ class EnvVars showBorders = GetEnvVar("SHOW_BORDERS" , 1); showIterations = GetEnvVar("SHOW_ITERATIONS" , 0); showPercentiles = GetEnvVarArray("SHOW_PERCENTILES", {}); + sweepMaxPow2 = GetEnvVar("SWEEP_MAX_POW2" , 29); + sweepMinPow2 = GetEnvVar("SWEEP_MIN_POW2" , 10); + sweepMinPow2 = std::clamp(sweepMinPow2, 0, 62); + sweepMaxPow2 = std::clamp(sweepMaxPow2, 0, 62); + if (sweepMinPow2 > sweepMaxPow2) std::swap(sweepMinPow2, sweepMaxPow2); + tdmBlockOrder = GetEnvVar("TDM_BLOCK_ORDER" , 0); + tdmBlockSize = GetEnvVar("TDM_BLOCK_SIZE" , 256); + tdmLdsBytes = GetEnvVar("TDM_LDS_BYTES" , 0); useHipEvents = GetEnvVar("USE_HIP_EVENTS" , 1); useHsaDma = GetEnvVar("USE_HSA_DMA" , 0); useInteractive = GetEnvVar("USE_INTERACTIVE" , 0); useSingleStream = GetEnvVar("USE_SINGLE_STREAM" , 1); validateDirect = GetEnvVar("VALIDATE_DIRECT" , 0); + validateOnDevice = GetEnvVar("VALIDATE_ON_DEVICE" , 0); validateSource = GetEnvVar("VALIDATE_SOURCE" , 0); ibGidIndex = GetEnvVar("IB_GID_INDEX" ,-1); @@ -390,28 +408,37 @@ 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(" SWEEP_MAX_POW2 - When 0 is specified for data size, this is the ending power of two exponent\n"); + printf(" SWEEP_MIN_POW2 - When 0 is specified for data size, this is the starting power of two exponent\n"); + printf(" TDM_BLOCK_ORDER - How blocks for TDM transfers are ordered. 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_LDS_BYTES - Amount of LDS bytes to allocate per workgroup for TDM kernels (0 = device max; K/M/G suffixes accepted)\n"); printf(" USE_HIP_EVENTS - Use HIP events for GFX executor timing\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"); printf(" VALIDATE_DIRECT - Validate GPU destination memory directly instead of staging GPU memory on host\n"); + printf(" VALIDATE_ON_DEVICE - Validate GPU dst/src memory via an on-device kernel instead of copying to host\n"); printf(" VALIDATE_SOURCE - Validate GPU src memory immediately after preparation\n"); printf("\n"); printf("Environment variables (back-end):\n"); printf("====================================\n"); - printf(" TB_RANK - Rank for socket communicator (0-based); defaults to 0 if unset or empty\n"); - printf(" TB_NUM_RANKS - Total ranks for socket mode (>=2); alone on rank 0 starts listener and logs worker env\n"); + printf(" TB_DUMP_CFG_FILE - Writes executed transfers to a config file\n"); + printf(" TB_DUMP_LINES - Dumps randomized input-line statistics for FILL_COMPRESS setup\n"); + printf(" TB_FORCE_SINGLE_POD - Forces all ranks into one pod (skips pod query)\n"); printf(" TB_MASTER_ADDR - Rank 0 hostname or IPv4 for workers; optional on rank 0 (auto-detected if unset)\n"); printf(" TB_MASTER_IFACE - When TB_MASTER_ADDR unset on rank 0, optional interface for IPv4 detection (e.g. eth0)\n"); printf(" TB_MASTER_PORT - Used to set Rank 0 port for socket communicator (default: 29500)\n"); + printf(" TB_NIC_FILTER - Regex filter to limit NIC visibility for NIC executors\n"); + printf(" TB_NUM_RANKS - Total ranks for socket mode (>=2); alone on rank 0 starts listener and logs worker env\n"); + printf(" TB_PAUSE - Pauses startup for debugger attachment\n"); + printf(" TB_RANK - Rank for socket communicator (0-based); defaults to 0 if unset or empty\n"); + printf(" TB_SEND_USLEEP - Microseconds to sleep after each socket SendData call (default: 0)\n"); + printf(" TB_SHOW_ALL_NUMA - Shows all CPU NUMA nodes incl. those with no cores (default: skip core-less)\n"); printf(" TB_SINGLE_LOG - In socket mode, only rank 0 logs when set\n"); printf(" TB_VERBOSE - Enables additional internal logging\n"); - printf(" TB_DUMP_CFG_FILE - Writes executed transfers to a config file\n"); - printf(" TB_DUMP_LINES - Dumps randomized input-line statistics for FILL_COMPRESS setup\n"); - printf(" TB_NIC_FILTER - Regex filter to limit NIC visibility for NIC executors\n"); - printf(" TB_FORCE_SINGLE_POD - Forces all ranks into one pod (skips pod query)\n"); printf(" TB_WALLCLOCK_RATE - Overrides queried GPU wallclock rate if needed\n"); - printf(" TB_PAUSE - Pauses startup for debugger attachment\n"); } void Print(std::string const& name, int32_t const value, const char* format, ...) const @@ -533,8 +560,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_ORDER", tdmBlockOrder, + "TDM Thread block ordering: %s", tdmBlockOrder == 0 ? "Sequential" : "Interleaved"); + Print("TDM_BLOCK_SIZE", tdmBlockSize, "TDM threadblock size of %d", tdmBlockSize); + Print("TDM_LDS_BYTES", tdmLdsBytes, "%s", + tdmLdsBytes == 0 ? "Using device max LDS bytes per workgroup" + : (std::string("Setting LDS to ") + std::to_string(tdmLdsBytes) + " bytes per workgroup").c_str()); 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, @@ -554,6 +587,8 @@ class EnvVars } Print("VALIDATE_DIRECT", validateDirect, "Validate GPU destination memory %s", validateDirect ? "directly" : "via CPU staging buffer"); + Print("VALIDATE_ON_DEVICE", validateOnDevice, + "Validate GPU memory %s", validateOnDevice ? "on-device via kernel" : "by copying to host"); Print("VALIDATE_SOURCE", validateSource, validateSource ? "Validate source after preparation" : "Do not perform source validation after prep"); printf("\n"); @@ -704,7 +739,9 @@ class EnvVars cfg.general.numSubIterations = numSubIterations; cfg.general.numWarmups = numWarmups; cfg.general.recordPerIteration = ((showIterations != 0) || !showPercentiles.empty()) ? 1 : 0; + cfg.general.useHipEvents = useHipEvents; cfg.general.useInteractive = useInteractive; + cfg.general.useMultiStream = !useSingleStream; cfg.data.alwaysValidate = alwaysValidate; cfg.data.blockBytes = blockBytes; @@ -712,9 +749,9 @@ class EnvVars cfg.data.fillCompress = fillCompress; cfg.data.fillPattern = fillPattern; cfg.data.validateDirect = validateDirect; + cfg.data.validateOnDevice = validateOnDevice; cfg.data.validateSource = validateSource; - cfg.dma.useHipEvents = useHipEvents; cfg.dma.useHsaCopy = useHsaDma; cfg.gfx.blockOrder = gfxBlockOrder; @@ -725,8 +762,6 @@ class EnvVars cfg.gfx.seType = gfxSeType; cfg.gfx.unrollFactor = gfxUnroll; cfg.gfx.temporalMode = gfxTemporal; - cfg.gfx.useHipEvents = useHipEvents; - cfg.gfx.useMultiStream = !useSingleStream; cfg.gfx.useSingleTeam = gfxSingleTeam; cfg.gfx.waveOrder = gfxWaveOrder; cfg.gfx.wordSize = gfxWordSize; @@ -741,6 +776,10 @@ class EnvVars cfg.nic.trafficClass = nicTrafficClass; cfg.nic.roceVersion = roceVersion; + cfg.tdm.blockOrder = tdmBlockOrder; + cfg.tdm.blockSize = tdmBlockSize; + cfg.tdm.ldsBytes = tdmLdsBytes; + return cfg; } }; diff --git a/src/client/Presets/AllToAll.hpp b/src/client/Presets/AllToAll.hpp index b9717a50..3db0709c 100644 --- a/src/client/Presets/AllToAll.hpp +++ b/src/client/Presets/AllToAll.hpp @@ -55,8 +55,19 @@ int AllToAllPreset(EnvVars& ev, int numSubExecs = EnvVars::GetEnvVar("NUM_SUB_EXEC" , 8); int showDetails = EnvVars::GetEnvVar("SHOW_DETAILS" , 0); int useDmaExec = EnvVars::GetEnvVar("USE_DMA_EXEC" , 0); + int useTdmExec = EnvVars::GetEnvVar("USE_TDM_EXEC" , 0); int useRemoteRead = EnvVars::GetEnvVar("USE_REMOTE_READ", 0); + // USE_DMA_EXEC and USE_TDM_EXEC are mutually exclusive; prefer DMA when both are requested + if (useDmaExec && useTdmExec) { + Utils::Print("[WARN] Both USE_DMA_EXEC and USE_TDM_EXEC are set. Using DMA executor\n"); + useTdmExec = 0; + } + + // Determine which GPU executor to use + ExeType exeType = useDmaExec ? EXE_GPU_DMA : (useTdmExec ? EXE_GPU_TDM : EXE_GPU_GFX); + char const* execName = useDmaExec ? "DMA" : (useTdmExec ? "TDM" : "GFX"); + // Check that all ranks have at least the number of GPUs requested // Warn if NIC configuration is slightly different from one another int numNics = TransferBench::GetNumExecutors(EXE_NIC, 0); @@ -108,6 +119,7 @@ int AllToAllPreset(EnvVars& ev, ev.Print("NUM_SUB_EXEC" , numSubExecs , "Using %d subexecutors/CUs per Transfer", numSubExecs); ev.Print("SHOW_DETAILS" , showDetails , "%s full Test details", showDetails ? "Showing" : "Hiding"); ev.Print("USE_DMA_EXEC" , useDmaExec , "Using %s executor", useDmaExec ? "DMA" : "GFX"); + ev.Print("USE_TDM_EXEC" , useTdmExec , "Using %s executor", useTdmExec ? "TDM" : "GFX"); ev.Print("USE_REMOTE_READ", useRemoteRead, "Using %s as executor", useRemoteRead ? "DST" : "SRC"); printf("\n"); } @@ -117,8 +129,8 @@ int AllToAllPreset(EnvVars& ev, Utils::Print("[ERROR] Cannot use %d GPUs. Detected %d GPUs\n", numGpus, numDetectedGpus); return ERR_FATAL; } - if (useDmaExec && (numSrcs != 1 || numDsts != 1)) { - Utils::Print("[ERROR] DMA execution can only be used for copies (A2A_MODE=0)\n"); + if ((useDmaExec || useTdmExec) && (numSrcs != 1 || numDsts != 1)) { + Utils::Print("[ERROR] %s execution can only be used for copies (A2A_MODE=0)\n", execName); return ERR_FATAL; } if (numResults * 2 > numRanks) { @@ -126,9 +138,6 @@ int AllToAllPreset(EnvVars& ev, return ERR_FATAL; } - // Collect the number of GPU devices to use - ExeType exeType = useDmaExec ? EXE_GPU_DMA : EXE_GPU_GFX; - std::vector, int>> reIndex(numRanks); std::vector transfers; for (int r = 0; r < numRanks; r++) { @@ -184,10 +193,10 @@ int AllToAllPreset(EnvVars& ev, } } - Utils::Print("GPU-%s All-To-All benchmark:\n", useDmaExec ? "DMA" : "GFX"); + Utils::Print("GPU-%s All-To-All benchmark:\n", execName); Utils::Print("==============================\n"); Utils::Print("[%lu bytes per Transfer] [%s:%d] [%d Read(s) %d Write(s)] [MemType:%s] [NIC QueuePairs:%d] [#Ranks:%d]\n", - numBytesPerTransfer, useDmaExec ? "DMA" : "GFX", numSubExecs, numSrcs, numDsts, + numBytesPerTransfer, execName, numSubExecs, numSrcs, numDsts, devMemTypeStr.c_str(), numQueuePairs, numRanks); if (transfers.size() == 0) { diff --git a/src/client/Presets/EmptyKernel.hpp b/src/client/Presets/EmptyKernel.hpp index c155059e..03b99937 100644 --- a/src/client/Presets/EmptyKernel.hpp +++ b/src/client/Presets/EmptyKernel.hpp @@ -101,10 +101,9 @@ int EmptyKernelPreset(EnvVars& ev, std::string const presetName, [[maybe_unused]] bool const bytesSpecified) { - if (Utils::GetNumRankGroups() > 1) { - Utils::Print("[ERROR] %s preset can only be run across ranks that are homogeneous\n", presetName.c_str()); + if (!Utils::AllRanksHaveSameGpuCount()) { + Utils::Print("[ERROR] %s preset requires all ranks to have the same number of GPUs\n", presetName.c_str()); Utils::Print("[ERROR] Run ./TransferBench without any args to display topology information\n"); - Utils::Print("[ERROR] TB_NIC_FILTER may also be used to limit NIC visibility\n"); return ERR_FATAL; } diff --git a/src/client/Presets/GfxSweep.hpp b/src/client/Presets/GfxSweep.hpp index 0816d82e..03079dfb 100644 --- a/src/client/Presets/GfxSweep.hpp +++ b/src/client/Presets/GfxSweep.hpp @@ -54,7 +54,6 @@ int GfxSweepPreset(EnvVars& ev, if (!ev.outputToCsv) Utils::Print("[GFX Sweep Related]\n"); ev.Print("BLOCKSIZES", blockList.size(), EnvVars::ToStr(blockList).c_str()); - ev.Print("GFX_TRANSFER", transferStr, "GFX Transfer to sweep (see config file format)"); ev.Print("KERNELS", kernelList.size(), EnvVars::ToStr(kernelList).c_str()); ev.Print("NUM_TRANSFERS", numTransfers, "Number of Transfers specified in GFX_TRANSFER"); ev.Print("NUM_SUB_EXECS", numSesList.size(), EnvVars::ToStr(numSesList).c_str()); @@ -63,6 +62,7 @@ int GfxSweepPreset(EnvVars& ev, ev.Print("UNROLLS", unrollList.size(), EnvVars::ToStr(unrollList).c_str()); ev.Print("WAVE_ORDERS", waveOrderList.size(), EnvVars::ToStr(waveOrderList).c_str()); ev.Print("WORDSIZES", wordSizeList.size(), EnvVars::ToStr(wordSizeList).c_str()); + ev.Print("GFX_TRANSFER", transferStr, "GFX Transfer to sweep (see config file format)"); Utils::Print("\n"); } } diff --git a/src/client/Presets/Help.hpp b/src/client/Presets/Help.hpp index 26ede846..052b707d 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("# Six 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) (Requires hardware support: AMD gfx1250 or NVIDIA sm_90+)\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,7 @@ 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: TDM-executor (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/Presets/PeerToPeer.hpp b/src/client/Presets/PeerToPeer.hpp index 5fbe1554..0032d854 100644 --- a/src/client/Presets/PeerToPeer.hpp +++ b/src/client/Presets/PeerToPeer.hpp @@ -35,6 +35,14 @@ int PeerToPeerPreset(EnvVars& ev, // Collect env vars for this preset int useDmaCopy = EnvVars::GetEnvVar("USE_GPU_DMA", 0); + int useTdmCopy = EnvVars::GetEnvVar("USE_GPU_TDM", 0); + + // USE_GPU_DMA and USE_GPU_TDM are mutually exclusive; prefer DMA when both are requested + if (useDmaCopy && useTdmCopy) { + Utils::Print("[WARN] Both USE_GPU_DMA and USE_GPU_TDM are set. Using DMA executor\n"); + useTdmCopy = 0; + } + char const* gpuExecName = useDmaCopy ? "DMA" : (useTdmCopy ? "TDM" : "GFX"); int cpuMemTypeIdx = EnvVars::GetEnvVar("CPU_MEM_TYPE", 0); int gpuMemTypeIdx = EnvVars::GetEnvVar("GPU_MEM_TYPE", 0); @@ -67,6 +75,7 @@ int PeerToPeerPreset(EnvVars& ev, ev.Print("SHOW_ITERATIONS", ev.showIterations, (ev.showIterations ? "Showing detailed iteration info" : "Showing compact info")); ev.Print("USE_GPU_DMA", useDmaCopy, "Using GPU-%s as GPU executor", useDmaCopy ? "DMA" : "GFX"); + ev.Print("USE_GPU_TDM", useTdmCopy, "Using GPU-%s as GPU executor", useTdmCopy ? "TDM" : "GFX"); ev.Print("USE_REMOTE_READ", useRemoteRead, "Using %s as executor", useRemoteRead ? "DST" : "SRC"); printf("\n"); } @@ -89,7 +98,7 @@ int PeerToPeerPreset(EnvVars& ev, printf("%sdirectional copy peak bandwidth GB/s [%s read / %s write] (GPU-Executor: %s)\n", isBidirectional ? "Bi" : "Uni", useRemoteRead ? "Remote" : "Local", useRemoteRead ? "Local" : "Remote", - useDmaCopy ? "DMA" : "GFX"); + gpuExecName); // Print header if (isBidirectional) { @@ -115,7 +124,7 @@ int PeerToPeerPreset(EnvVars& ev, double avgBwSum[2][2] = {}; int avgCount[2][2] = {}; - ExeType const gpuExeType = useDmaCopy ? EXE_GPU_DMA : EXE_GPU_GFX; + ExeType const gpuExeType = useDmaCopy ? EXE_GPU_DMA : (useTdmCopy ? EXE_GPU_TDM : EXE_GPU_GFX); // Loop over all possible src/dst pairs for (int src = 0; src < numDevices; src++) { diff --git a/src/client/Presets/PodAllToAll.hpp b/src/client/Presets/PodAllToAll.hpp index 2b3d4f62..bfc29948 100644 --- a/src/client/Presets/PodAllToAll.hpp +++ b/src/client/Presets/PodAllToAll.hpp @@ -64,10 +64,21 @@ int PodAllToAllPreset(EnvVars& ev, int numSubExecs = EnvVars::GetEnvVar("NUM_SUB_EXEC" , 8); int showDetails = EnvVars::GetEnvVar("SHOW_DETAILS" , 0); int useDmaExec = EnvVars::GetEnvVar("USE_DMA_EXEC" , 0); + int useTdmExec = EnvVars::GetEnvVar("USE_TDM_EXEC" , 0); int useRemoteRead = EnvVars::GetEnvVar("USE_REMOTE_READ", 0); int groupStride = EnvVars::GetEnvVar("GROUP_STRIDE" , 1); int numGroups = EnvVars::GetEnvVar("NUM_GROUPS" , 1); + // USE_DMA_EXEC and USE_TDM_EXEC are mutually exclusive; prefer DMA when both are requested + if (useDmaExec && useTdmExec) { + Utils::Print("[WARN] Both USE_DMA_EXEC and USE_TDM_EXEC are set. Using DMA executor\n"); + useTdmExec = 0; + } + + // Determine which GPU executor to use + ExeType exeType = useDmaExec ? EXE_GPU_DMA : (useTdmExec ? EXE_GPU_TDM : EXE_GPU_GFX); + char const* execName = useDmaExec ? "DMA" : (useTdmExec ? "TDM" : "GFX"); + // Check that all ranks have at least the number of GPUs requested // Warn if NIC configuration is slightly different from one another int numNics = TransferBench::GetNumExecutors(EXE_NIC, 0); @@ -115,6 +126,7 @@ int PodAllToAllPreset(EnvVars& ev, ev.Print("NUM_QUEUE_PAIRS", numQueuePairs, "Using %d queue pairs for NIC transfers", numQueuePairs); ev.Print("NUM_SUB_EXEC" , numSubExecs , "Using %d subexecutors/CUs per Transfer", numSubExecs); ev.Print("USE_DMA_EXEC" , useDmaExec , "Using %s executor", useDmaExec ? "DMA" : "GFX"); + ev.Print("USE_TDM_EXEC" , useTdmExec , "Using %s executor", useTdmExec ? "TDM" : "GFX"); ev.Print("USE_REMOTE_READ", useRemoteRead, "Using %s as executor", useRemoteRead ? "DST" : "SRC"); ev.Print("GROUP_STRIDE" , groupStride , "Stride permutation on device list before splitting into groups"); ev.Print("NUM_GROUPS" , numGroups , "Splitting each pod into %d group(s) for a2a", numGroups); @@ -126,8 +138,8 @@ int PodAllToAllPreset(EnvVars& ev, Utils::Print("[ERROR] Cannot use %d GPUs. Detected %d GPUs\n", numGpus, numDetectedGpus); return ERR_FATAL; } - if (useDmaExec && (numSrcs != 1 || numDsts != 1)) { - Utils::Print("[ERROR] DMA execution can only be used for copies (A2A_MODE=0)\n"); + if ((useDmaExec || useTdmExec) && (numSrcs != 1 || numDsts != 1)) { + Utils::Print("[ERROR] %s execution can only be used for copies (A2A_MODE=0)\n", execName); return ERR_FATAL; } @@ -136,14 +148,13 @@ int PodAllToAllPreset(EnvVars& ev, return ERR_FATAL; } - Utils::Print("GPU-%s IntraPod All-To-All benchmark:\n", useDmaExec ? "DMA" : "GFX"); + Utils::Print("GPU-%s IntraPod All-To-All benchmark:\n", execName); Utils::Print("==============================\n"); Utils::Print("[%lu bytes per Transfer] [%s:%d] [%d Read(s) %d Write(s)] [MemType:%s] [NIC QueuePairs:%d] [#Ranks:%d]\n", - numBytesPerTransfer, useDmaExec ? "DMA" : "GFX", numSubExecs, numSrcs, numDsts, + numBytesPerTransfer, execName, numSubExecs, numSrcs, numDsts, devMemTypeStr.c_str(), numQueuePairs, numRanks); TransferBench::ConfigOptions cfg = ev.ToConfigOptions(); - ExeType exeType = useDmaExec ? EXE_GPU_DMA : EXE_GPU_GFX; Utils::RankPerPodMap& rankToPod = Utils::GetRankPerPodMap(); for (auto const& [pod, ranks] : rankToPod) { diff --git a/src/client/Presets/PodPeerToPeer.hpp b/src/client/Presets/PodPeerToPeer.hpp index fe1cc775..2abcd308 100644 --- a/src/client/Presets/PodPeerToPeer.hpp +++ b/src/client/Presets/PodPeerToPeer.hpp @@ -25,9 +25,8 @@ int PodPeerToPeerPreset(EnvVars& ev, std::string const presetName, bool const bytesSpecified) { - if (Utils::GetNumRankGroups() > 1) { - Utils::Print("[ERROR] Pod p2p preset can only be run across ranks that are homogenous\n"); - Utils::Print("[ERROR] All ranks currently have to be under the same physical and virtual pod\n"); + if (!Utils::AllRanksHaveSameGpuCount()) { + Utils::Print("[ERROR] Pod p2p preset requires all ranks to have the same number of GPUs\n"); Utils::Print("[ERROR] Run ./TransferBench without any args to display topology information\n"); return ERR_FATAL; } @@ -42,6 +41,14 @@ int PodPeerToPeerPreset(EnvVars& ev, // Collect env vars for this preset int useDmaCopy = EnvVars::GetEnvVar("USE_GPU_DMA", 0); + int useTdmCopy = EnvVars::GetEnvVar("USE_GPU_TDM", 0); + + // USE_GPU_DMA and USE_GPU_TDM are mutually exclusive; prefer DMA when both are requested + if (useDmaCopy && useTdmCopy) { + Utils::Print("[WARN] Both USE_GPU_DMA and USE_GPU_TDM are set. Using DMA executor\n"); + useTdmCopy = 0; + } + char const* gpuExecName = useDmaCopy ? "DMA" : (useTdmCopy ? "TDM" : "GFX"); int gpuMemTypeIdx = EnvVars::GetEnvVar("GPU_MEM_TYPE", 0); int numGpuDevices = EnvVars::GetEnvVar("NUM_GPU_DEVICES", numDetectedGpus); int numGpuSubExecs = EnvVars::GetEnvVar("NUM_GPU_SE", useDmaCopy ? 1 : TransferBench::GetNumSubExecutors({EXE_GPU_GFX, 0})); @@ -67,6 +74,7 @@ int PodPeerToPeerPreset(EnvVars& ev, : "Bidirectional"); ev.Print("PARALLEL_LVL", parallelLevel, "Executing p2p in parallel level %d (0: no parallel, 1: node pairs in parallel)", parallelLevel); ev.Print("USE_GPU_DMA", useDmaCopy, "Using GPU-%s as GPU executor", useDmaCopy ? "DMA" : "GFX"); + ev.Print("USE_GPU_TDM", useTdmCopy, "Using GPU-%s as GPU executor", useTdmCopy ? "TDM" : "GFX"); ev.Print("USE_REMOTE_READ", useRemoteRead, "Using %s as executor", useRemoteRead ? "DST" : "SRC"); printf("\n"); } @@ -95,7 +103,7 @@ int PodPeerToPeerPreset(EnvVars& ev, for (int i = 0; i < n; i++) deviceLookup[{devices[i].memRank, devices[i].memIndex}] = i; - ExeType const gpuExeType = useDmaCopy ? EXE_GPU_DMA : EXE_GPU_GFX; + ExeType const gpuExeType = useDmaCopy ? EXE_GPU_DMA : (useTdmCopy ? EXE_GPU_TDM : EXE_GPU_GFX); for (int isBidirectional = 0; isBidirectional <= 1; isBidirectional++) { if ((p2pMode == 1 && isBidirectional == 1) || @@ -105,7 +113,7 @@ int PodPeerToPeerPreset(EnvVars& ev, isBidirectional ? "Bi" : "Uni", useRemoteRead ? "Remote" : "Local", useRemoteRead ? "Local" : "Remote", - useDmaCopy ? "DMA" : "GFX"); + gpuExecName); std::vector avgBandwidth(n * n, 0.0); diff --git a/src/client/Presets/Presets.hpp b/src/client/Presets/Presets.hpp index 4bd631f7..8465e389 100644 --- a/src/client/Presets/Presets.hpp +++ b/src/client/Presets/Presets.hpp @@ -50,6 +50,7 @@ THE SOFTWARE. #include "Schmoo.hpp" #include "SmokeTest.hpp" #include "Sweep.hpp" +#include "TdmSweep.hpp" #include "WallClock.hpp" typedef int (*PresetFunc)(EnvVars& ev, @@ -88,6 +89,7 @@ std::map presetFuncMap = {"schmoo", {SchmooPreset, "Scaling tests for local/remote read/write/copy"}}, {"smoketest", {SmokeTestPreset, "Simple correctness smoke-test"}}, {"sweep", {SweepPreset, "Ordered sweep through sets of Transfers"}}, + {"tdmsweep", {TdmSweepPreset, "Sweep over TDM executor options (block size / LDS / order / subExecs) for a given TDM Transfer"}}, {"wallclock", {WallClockPreset, "Tests wallclock consistency across XCCs within a GPU"}}, }; diff --git a/src/client/Presets/Rings.hpp b/src/client/Presets/Rings.hpp index 90914159..9226370f 100644 --- a/src/client/Presets/Rings.hpp +++ b/src/client/Presets/Rings.hpp @@ -28,11 +28,10 @@ int RingsPreset(EnvVars& ev, std::string const presetName, bool const bytesSpecified) { - // Check for homogeneous ranks - if (Utils::GetNumRankGroups() > 1) { - Utils::Print("[ERROR] rings preset can only be run across ranks that are homogeneous\n"); + // Check that all ranks have the same number of GPUs + if (!Utils::AllRanksHaveSameGpuCount()) { + Utils::Print("[ERROR] rings preset requires all ranks to have the same number of GPUs\n"); Utils::Print("[ERROR] Run ./TransferBench without any args to display topology information\n"); - Utils::Print("[ERROR] TB_NIC_FILTER may also be used to limit NIC visibility\n"); return ERR_FATAL; } @@ -51,10 +50,21 @@ int RingsPreset(EnvVars& ev, int numSubExecs = EnvVars::GetEnvVar("NUM_SUB_EXEC" , 8); int showDetails = EnvVars::GetEnvVar("SHOW_DETAILS" , 0); int useDmaExec = EnvVars::GetEnvVar("USE_DMA_EXEC" , 0); + int useTdmExec = EnvVars::GetEnvVar("USE_TDM_EXEC" , 0); int useRemoteRead = EnvVars::GetEnvVar("USE_REMOTE_READ", 0); int stride = EnvVars::GetEnvVar("STRIDE" , 1); int ringSize = EnvVars::GetEnvVar("RING_SIZE" , numRanks * numGpus); + // USE_DMA_EXEC and USE_TDM_EXEC are mutually exclusive; prefer DMA when both are requested + if (useDmaExec && useTdmExec) { + Utils::Print("[WARN] Both USE_DMA_EXEC and USE_TDM_EXEC are set. Using DMA executor\n"); + useTdmExec = 0; + } + + // Determine which GPU executor to use + ExeType exeType = useDmaExec ? EXE_GPU_DMA : (useTdmExec ? EXE_GPU_TDM : EXE_GPU_GFX); + char const* execName = useDmaExec ? "DMA" : (useTdmExec ? "TDM" : "GFX"); + if (numGpus <= 0 || numGpus > numDetectedGpus) { Utils::Print("[ERROR] Cannot use %d GPUs. Detected %d GPUs\n", numGpus, numDetectedGpus); @@ -87,6 +97,7 @@ int RingsPreset(EnvVars& ev, ev.Print("NUM_QUEUE_PAIRS", numQueuePairs, "Using %d queue pairs for NIC transfers", numQueuePairs); ev.Print("NUM_SUB_EXEC" , numSubExecs , "Using %d subexecutors/CUs per Transfer", numSubExecs); ev.Print("USE_DMA_EXEC" , useDmaExec , "Using %s executor", useDmaExec ? "DMA" : "GFX"); + ev.Print("USE_TDM_EXEC" , useTdmExec , "Using %s executor", useTdmExec ? "TDM" : "GFX"); ev.Print("USE_REMOTE_READ", useRemoteRead, "Using %s as executor", useRemoteRead ? "DST" : "SRC"); ev.Print("STRIDE" , stride , "Reordering devices by taking %d steps", stride); ev.Print("RING_SIZE" , ringSize , "Building rings of size %d", ringSize); @@ -94,14 +105,13 @@ int RingsPreset(EnvVars& ev, } } - Utils::Print("GPU-%s Rings benchmark:\n", useDmaExec ? "DMA" : "GFX"); + Utils::Print("GPU-%s Rings benchmark:\n", execName); Utils::Print("==============================\n"); Utils::Print("[%lu bytes per Transfer] [%s:%d] [MemType:%s] [NIC QueuePairs:%d] [#Ranks:%d]\n", - numBytesPerTransfer, useDmaExec ? "DMA" : "GFX", numSubExecs, + numBytesPerTransfer, execName, numSubExecs, devMemTypeStr.c_str(), numQueuePairs, numRanks); TransferBench::ConfigOptions cfg = ev.ToConfigOptions(); - ExeType exeType = useDmaExec ? EXE_GPU_DMA : EXE_GPU_GFX; int numRings = totalGpus / ringSize; Utils::Print("Running %d parallel ring(s) each of %d devices. All numbers in GB/s:\n", numRings, ringSize); @@ -192,7 +202,7 @@ int RingsPreset(EnvVars& ev, for (int i = 0; i < colsPerRing; i++) table.Set(headerRow, currCol+i, "Ring%02d", ringIdx); table.Set(headerRow+1, currCol, "Device"); - table.Set(headerRow+1, currCol+1, "%s BW", useDmaExec ? "DMA" : "GFX"); + table.Set(headerRow+1, currCol+1, "%s BW", execName); if (numQueuePairs) { table.Set(headerRow+1, currCol+2, "NIC BW"); } diff --git a/src/client/Presets/SmokeTest.hpp b/src/client/Presets/SmokeTest.hpp index 4b3325b8..b6880f91 100644 --- a/src/client/Presets/SmokeTest.hpp +++ b/src/client/Presets/SmokeTest.hpp @@ -21,6 +21,8 @@ THE SOFTWARE. */ #include +#include +#include namespace { @@ -43,7 +45,8 @@ int RunTest(int testNum, bool isParallel, bool useBdma, int targetGpu, - int totalGpus) + int totalGpus, + std::vector& failLog) { int numFail = 0; @@ -114,10 +117,10 @@ int RunTest(int testNum, allTransfers.clear(); // Combine transfers from each GPU and run them all in parallel (unless isParallel=false) - for (int rank = 0; allPass && rank < numRanks; rank++) { + for (int rank = 0; rank < numRanks; rank++) { if (!isParallel && rank != targetRank) continue; int numGpus = GetNumExecutors(exeType, rank); - for (int gpuIdx = 0; allPass && gpuIdx < numGpus; gpuIdx++) { + for (int gpuIdx = 0; gpuIdx < numGpus; gpuIdx++) { if (!isParallel && gpuIdx != targetIdx) continue; if (isH2D || isD2H) { // Copy to/from closest CPU NUMA node for this GPU @@ -165,7 +168,10 @@ int RunTest(int testNum, if (isBroadcast || isGather) { if (!RunTransfers(cfg, transfers, results)) { allPass = false; - break; + std::string entry = "Test " + std::to_string(testNum) + " | " + std::string(transferStr) + " | numBytes=" + std::to_string(numBytes); + for (auto const& e : results.errResults) + if (e.errType == ERR_FATAL) entry += "\n ERROR: " + e.errMsg; + failLog.push_back(entry); } } else { // Otherwise accumulate the transfers to run in parallel allTransfers.insert(allTransfers.end(), transfers.begin(), transfers.end()); @@ -175,6 +181,19 @@ int RunTest(int testNum, if (!(isBroadcast || isGather)) { if (!RunTransfers(cfg, allTransfers, results)) { allPass = false; + // Build a summary entry covering all accumulated transfers + std::string entry = "Test " + std::to_string(testNum) + " | numBytes=" + std::to_string(numBytes) + " | transfers:"; + for (auto const& t : allTransfers) { + char buf[MAX_TRANSFER_STRLEN]; + snprintf(buf, MAX_TRANSFER_STRLEN, " (R%d%c%d->R%d%c%d->R%d%c%d)", + t.srcs.empty() ? 0 : t.srcs[0].memRank, t.srcs.empty() ? '?' : MemTypeStr[t.srcs[0].memType], t.srcs.empty() ? 0 : t.srcs[0].memIndex, + t.exeDevice.exeRank, ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex, + t.dsts.empty() ? 0 : t.dsts[0].memRank, t.dsts.empty() ? '?' : MemTypeStr[t.dsts[0].memType], t.dsts.empty() ? 0 : t.dsts[0].memIndex); + entry += buf; + } + for (auto const& e : results.errResults) + if (e.errType == ERR_FATAL) entry += "\n ERROR: " + e.errMsg; + failLog.push_back(entry); } } Utils::Print("%s", allPass ? pass.c_str() : fail.c_str()); fflush(stdout); @@ -310,6 +329,7 @@ int SmokeTestPreset(EnvVars& ev, std::string l2(lPad2Size, ' '), r2(rPad2Size, ' '); int testsFailed = 0; + std::vector failLog; auto ExecuteTests = [&](std::string label, int x, int y, bool isParallel) { int numLines = isParallel ? 1 : totalGpus; @@ -328,7 +348,7 @@ int SmokeTestPreset(EnvVars& ev, Utils::Print("%s", l2.c_str()); fflush(stdout); - testsFailed += RunTest(x, testsToRun, sizeList, 1, cfg, cpuMemType, gpuMemType, seMaxBytes, isParallel, useBdma, line, totalGpus); + testsFailed += RunTest(x, testsToRun, sizeList, 1, cfg, cpuMemType, gpuMemType, seMaxBytes, isParallel, useBdma, line, totalGpus, failLog); Utils::Print("%s|", r2.c_str()); if (line == 0) { Utils::Print(" %02d |", y); @@ -338,7 +358,7 @@ int SmokeTestPreset(EnvVars& ev, for (auto numSubExec : gfxSesList) { Utils::Print("%s", l2.c_str()); fflush(stdout); - testsFailed += RunTest(y, testsToRun, sizeList, numSubExec, cfg, cpuMemType, gpuMemType, seMaxBytes, isParallel, useBdma, line, totalGpus); + testsFailed += RunTest(y, testsToRun, sizeList, numSubExec, cfg, cpuMemType, gpuMemType, seMaxBytes, isParallel, useBdma, line, totalGpus, failLog); Utils::Print("%s|", r2.c_str()); } Utils::Print("\n"); @@ -379,7 +399,21 @@ int SmokeTestPreset(EnvVars& ev, // Show summary if (testsFailed) { - Utils::Print("[WARN] %d Tests FAILED\n", testsFailed); + // Write failure details to a log file + char logPath[128]; + snprintf(logPath, sizeof(logPath), "/tmp/transferbench_smoketest_%d.log", (int)getpid()); + FILE* logFile = fopen(logPath, "w"); + if (logFile) { + time_t now = time(nullptr); + fprintf(logFile, "TransferBench smoketest failures — %s\n", ctime(&now)); + fprintf(logFile, "Total failed tests: %d\n\n", testsFailed); + for (size_t i = 0; i < failLog.size(); i++) + fprintf(logFile, "[%zu] %s\n\n", i + 1, failLog[i].c_str()); + fclose(logFile); + Utils::Print("[WARN] %d Tests FAILED — details written to %s\n", testsFailed, logPath); + } else { + Utils::Print("[WARN] %d Tests FAILED (could not write log to %s)\n", testsFailed, logPath); + } } else { Utils::Print("All tests passed\n"); } diff --git a/src/client/Presets/TdmSweep.hpp b/src/client/Presets/TdmSweep.hpp new file mode 100644 index 00000000..31b3845f --- /dev/null +++ b/src/client/Presets/TdmSweep.hpp @@ -0,0 +1,263 @@ +/* +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. +*/ + +#include "EnvVars.hpp" + +// TdmSweepPreset - sweeps every knob that affects Tensor-Data-Mover (TDM) copy +// performance for a given TDM Transfer and reports the best-performing +// combination, mirroring the structure of the "gfxsweep" preset. +// +// The TDM executor (EXE_GPU_TDM, letter 'T' in a Transfer expr) stages HBM->LDS +// ->HBM without touching cache. Its performance is governed by: +// * threadblock size (TDM_BLOCK_SIZE : # threads/block, multiple of 32) +// * LDS staging window (TDM_LDS_BYTES : bytes/block, 0 = device max) +// * threadblock order (TDM_BLOCK_ORDER: 0=sequential 1=interleaved 2=random) +// * number of subExecs (threadblocks/WGPs engaged, NUM_SUB_EXECS) +// This preset sweeps the cartesian product of all four and finds the peak. +int TdmSweepPreset(EnvVars& ev, + size_t const numBytesPerTransfer, + std::string const presetName, + bool const bytesSpecified) +{ + enum TimingMode + { + TimingModeAuto = -1, + TimingModeCpu = 0, + TimingModeHip = 1, + TimingModeGpu = 2 + }; + + // Verify the hardware can actually run TDM copies before sweeping anything. + int const numTdmGpus = TransferBench::GetNumExecutors(EXE_GPU_TDM); + if (numTdmGpus <= 0 || !tdm::IsTdmCopySupported(0)) { + Utils::Print("[WARN] TDM executor is not supported on this device " + "(requires TDM-capable hardware: gfx1250 or NVIDIA sm_90+). Terminating tdmsweep.\n"); + return ERR_FATAL; + } + + // Collect environment variables for this preset + vector blockList = EnvVars::GetEnvVarArray("BLOCKSIZES", {64,128,256,512,1024}); + vector blockOrders = EnvVars::GetEnvVarArray("BLOCK_ORDERS", {0}); + vector ldsList = EnvVars::GetEnvVarArray("LDS_BYTES", {8192,16384,32768,65536,0}); + vector numSesList = EnvVars::GetEnvVarArray("NUM_SUB_EXECS", {2,4,8,16,32,64}); + int numTransfers = EnvVars::GetEnvVar( "NUM_TRANSFERS", 1); + int timingMode = EnvVars::GetEnvVar( "TIMING_MODE", TimingModeAuto); + std::string transferStr = EnvVars::GetEnvVar( "TDM_TRANSFER", "G0->T0->G0"); + + // Print off relevant environment variables + if (Utils::RankDoesOutput()) { + if (!ev.hideEnv) { + ev.DisplayEnvVars(); + if (!ev.outputToCsv) + Utils::Print("[TDM Sweep Related]\n"); + ev.Print("BLOCKSIZES", blockList.size(), EnvVars::ToStr(blockList).c_str()); + ev.Print("BLOCK_ORDERS", blockOrders.size(), "%s (0=sequential 1=interleaved 2=random)", EnvVars::ToStr(blockOrders).c_str()); + ev.Print("LDS_BYTES", ldsList.size(), "%s (0 = device max LDS per block)", EnvVars::ToStr(ldsList).c_str()); + ev.Print("NUM_SUB_EXECS", numSesList.size(), EnvVars::ToStr(numSesList).c_str()); + ev.Print("NUM_TRANSFERS", numTransfers, "Number of Transfers specified in TDM_TRANSFER"); + ev.Print("TIMING_MODE", timingMode, "-1=auto 0=Aggregate CPU, 1=Executor Time, 2=Transfer Time"); + ev.Print("TDM_TRANSFER", transferStr, "TDM Transfer to sweep (see config file format)"); + Utils::Print("\n"); + } + } + + if (timingMode < TimingModeAuto || timingMode > TimingModeGpu) { + Utils::Print("TIMING_MODE value is invalid (%d)\n", timingMode); + return ERR_FATAL; + } + + if (numSesList.empty()) { + Utils::Print("NUM_SUB_EXECS should not be empty\n"); + return ERR_FATAL; + } + + // TDM block size must be a positive multiple of 32 + for (int bs : blockList) { + if (bs <= 0 || bs % 32) { + Utils::Print("[ERROR] BLOCKSIZES value %d is invalid (TDM block size must be a positive multiple of 32)\n", bs); + return ERR_FATAL; + } + } + for (int lds : ldsList) { + if (lds < 0) { + Utils::Print("[ERROR] LDS_BYTES value %d is invalid (must be >= 0; 0 = device max)\n", lds); + return ERR_FATAL; + } + } + for (int bo : blockOrders) { + if (bo < 0 || bo > 2) { + Utils::Print("[ERROR] BLOCK_ORDERS value %d is invalid (must be 0, 1, or 2)\n", bo); + return ERR_FATAL; + } + } + + std::vector transfers; + Utils::CheckForError(ParseTransfers(std::to_string(numTransfers) + " 1 " + transferStr, transfers)); + if (transfers.size() == 0) { + Utils::Print("[WARN] No valid Transfers found in TDM_TRANSFER\n"); + return 0; + } + + // Automatically pick timing method + if (timingMode == TimingModeAuto) { + // Use Transfer timing if there is only one Transfer + if (transfers.size() == 1) timingMode = TimingModeGpu; + // Use Executor timing if there is only one executor + else { + bool singleExecutor = true; + for (size_t i = 1; i < transfers.size(); i++) { + if (transfers[i].exeDevice < transfers[0].exeDevice || + transfers[0].exeDevice < transfers[i].exeDevice || + transfers[i].exeSubIndex != transfers[0].exeSubIndex || + transfers[i].exeSubSlot != transfers[0].exeSubSlot) { + singleExecutor = false; + break; + } + } + timingMode = singleExecutor ? TimingModeHip : TimingModeCpu; + } + } + if (timingMode < 0 || timingMode > 2) { + Utils::Print("[ERROR] Invalid timing mode %d\n", timingMode); + return ERR_FATAL; + } + + // Print out the Transfers being run + Utils::Print("TDM sweep: (%lu bytes per Transfer). All values are %s-timed GB/s\n", numBytesPerTransfer, + timingMode == TimingModeCpu ? "Aggregate-CPU" : + timingMode == TimingModeHip ? "HIP-event" : + "GPU wallclock"); + Utils::Print("=======================================================================================\n"); + + bool isMultiNode = GetNumRanks() > 1; + for (size_t i = 0; i < transfers.size(); i++) { + Transfer& t = transfers[i]; + Utils::Print("Transfer %5lu: (%s->", i, Utils::MemDevicesToStr(t.srcs).c_str()); + if (isMultiNode) Utils::Print("R%d", t.exeDevice.exeRank); + Utils::Print("%c%d", ExeTypeStr[t.exeDevice.exeType], t.exeDevice.exeIndex); + if (t.exeDevice.exeSlot) Utils::Print("%c", 'A' + t.exeDevice.exeSlot); + if (t.exeSubIndex != -1) Utils::Print(".%d", t.exeSubIndex); + if (t.exeSubSlot != 0) Utils::Print("%c", 'A' + t.exeSubSlot); + Utils::Print("->%s)\n", Utils::MemDevicesToStr(t.dsts).c_str()); + + if (t.exeDevice.exeType != EXE_GPU_TDM) { + Utils::Print("[ERROR] tdmsweep preset only works on Transfers that are using the TDM executor " + "(use the 'T' executor, e.g. TDM_TRANSFER=\"G0->T0->G1\")\n"); + return ERR_FATAL; + } + t.numBytes = numBytesPerTransfer; + } + + Utils::Print("=======================================================================================\n"); + + ConfigOptions cfg = ev.ToConfigOptions(); + + // Print header + char sep = ev.outputToCsv ? ',' : ' '; + Utils::Print(" BlkO %c BlkS %c LDSBytes ", sep, sep); + for (int numSubExec : numSesList) + Utils::Print("%c SE %03d", sep, numSubExec); + Utils::Print("\n"); + + int bestSe = -1; + double overallBestBw = 0; + vector bestBw(numSesList.size(), 0.0); + // best[s] = {blockOrder, blockSize, ldsBytes, numSubExec} + vector> best(numSesList.size(), vector(4)); + + // Loop over all combinations + for (int blockOrder : blockOrders) { cfg.tdm.blockOrder = blockOrder; + for (int blockSize : blockList) { cfg.tdm.blockSize = blockSize; + for (int ldsBytes : ldsList) { cfg.tdm.ldsBytes = ldsBytes; + Utils::Print(" %1d %c %4d %c %8d ", blockOrder, sep, blockSize, sep, ldsBytes); + fflush(stdout); + for (auto s = 0; s < numSesList.size(); s++) { + int numSubExec = numSesList[s]; + for (Transfer& t : transfers) t.numSubExecs = numSubExec; + + TestResults result; + // A given combination may be rejected by the library (e.g. LDS window + // larger than the device max). Treat that as a skipped cell (N/A) and + // keep sweeping instead of aborting the whole matrix. + if (RunTransfers(cfg, transfers, result)) { + double bw = 0.0; + switch (timingMode) { + case 0: bw = result.avgTotalBandwidthGbPerSec; break; + case 1: + for (auto const& e : result.exeResults) { + bw = std::max(bw, e.second.avgBandwidthGbPerSec); + } + break; + case 2: default: + for (auto const& t : result.tfrResults) { + bw = std::max(bw, t.avgBandwidthGbPerSec); + } + break; + } + + if (bw > bestBw[s]) { + bestBw[s] = bw; + best[s] = {blockOrder, blockSize, ldsBytes, numSubExec}; + if (bw > overallBestBw) { + overallBestBw = bw; + bestSe = s; + } + } + Utils::Print("%c%8.2f", sep, bw); + } else { + Utils::Print("%c%8s", sep, "N/A"); + } + fflush(stdout); + } + Utils::Print("\n"); + fflush(stdout); + } + } + } + + Utils::Print(" BlkO %c BlkS %c LDSBytes ", sep, sep); + for (auto s = 0; s < numSesList.size(); s++) { + Utils::Print("%c%8.2f", sep, bestBw[s]); + } + Utils::Print("\n"); + + if (bestSe == -1) { + Utils::Print("[WARN] No transfers executed successfully - check sweep parameters and TDM support\n"); + return ERR_FATAL; + } + + // Print combination that produced highest bandwidth + Utils::Print("=======================================================================================\n"); + Utils::Print("Highest bandwidth found: %7.2f GB/s (%s-timed)\n", overallBestBw, + timingMode == TimingModeCpu ? "Aggregate-CPU" : + timingMode == TimingModeHip ? "HIP-event" : + "GPU wallclock"); + Utils::Print(" BlockOrder : %7d [TDM_BLOCK_ORDER=%d]\n", best[bestSe][0], best[bestSe][0]); + Utils::Print(" BlockSize : %7d [TDM_BLOCK_SIZE=%d]\n", best[bestSe][1], best[bestSe][1]); + Utils::Print(" LDS Bytes : %7d [TDM_LDS_BYTES=%d]\n", best[bestSe][2], best[bestSe][2]); + Utils::Print(" NumSubExec : %7d\n", best[bestSe][3]); + Utils::Print("Command to run best result:\n"); + Utils::Print("TDM_BLOCK_ORDER=%d TDM_BLOCK_SIZE=%d TDM_LDS_BYTES=%d ./TransferBench cmdline %lu \"%d %d %s\"\n", + best[bestSe][0], best[bestSe][1], best[bestSe][2], + numBytesPerTransfer, numTransfers, best[bestSe][3], transferStr.c_str()); + return ERR_NONE; +} diff --git a/src/client/Presets/WallClock.hpp b/src/client/Presets/WallClock.hpp index 4b441dff..6f4f5b65 100644 --- a/src/client/Presets/WallClock.hpp +++ b/src/client/Presets/WallClock.hpp @@ -54,8 +54,7 @@ __global__ void GetTimestamps(uint64_t* timestamps, auto start = GetTimestamp(); // Collect XCD for this - int xccId; - GetXccId(xccId); + uint32_t xccId = GetXccId(); int idx = (indexType == 0) ? xccId : blockIdx.x; if (xccMask & (1U << xccId)) { timestamps[idx] = 0; @@ -99,11 +98,10 @@ int WallClockPreset(EnvVars& ev, int numRanks = GetNumRanks(); int myRank = GetRank(); - // Check for single homogenous group - if (Utils::GetNumRankGroups() > 1) { - Utils::Print("[ERROR] wallclock preset can only be run across ranks that are homogenous\n"); + // Check that all ranks have the same number of GPUs + if (!Utils::AllRanksHaveSameGpuCount()) { + Utils::Print("[ERROR] wallclock preset requires all ranks to have the same number of GPUs\n"); Utils::Print("[ERROR] Run ./TransferBench without any args to display topology information\n"); - Utils::Print("[ERROR] TB_NIC_FILTER may also be used to limit NIC visibility\n"); return ERR_FATAL; } diff --git a/src/client/Topology.hpp b/src/client/Topology.hpp index 1f5501eb..98cb717f 100644 --- a/src/client/Topology.hpp +++ b/src/client/Topology.hpp @@ -27,20 +27,9 @@ THE SOFTWARE. static int RemappedCpuIndex(int origIdx) { - static std::vector remappingCpu; - - // Build CPU remapping on first use - // Skip numa nodes that are not configured - if (remappingCpu.empty()) { - for (int node = 0; node <= numa_max_node(); node++) { - if (numa_bitmask_isbitset(numa_get_mems_allowed(), node)) - remappingCpu.push_back(node); - else { - remappingCpu.push_back(-1); - } - } - } - return remappingCpu[origIdx]; + // Map a logical CPU NUMA index to its physical NUMA node using the same mapping + // the core library uses (honors TB_SHOW_ALL_NUMA and core-less node skipping). + return TransferBench::GetCpuNumaPhysicalNode(origIdx); } static void PrintNicToGPUTopo(bool outputToCsv) @@ -52,16 +41,21 @@ static void PrintNicToGPUTopo(bool outputToCsv) int numGpus = TransferBench::GetNumExecutors(EXE_GPU_GFX); auto const& ibvDeviceList = GetIbvDeviceList(); - for (int i = 0; i < ibvDeviceList.size(); i++) { - std::string closestGpusStr = ""; - for (int j = 0; j < numGpus; j++) { - if (TransferBench::GetClosestNicToGpu(j) == i) { - if (closestGpusStr != "") closestGpusStr += ","; - closestGpusStr += std::to_string(j); - } + // Build inverse map: for each NIC, which GPUs list it as closest? + std::vector closestGpusForNic(ibvDeviceList.size(), ""); + for (int j = 0; j < numGpus; j++) { + std::vector nicsForGpu; + TransferBench::GetClosestNicsToGpu(nicsForGpu, j); + for (int nicIdx : nicsForGpu) { + if (nicIdx < 0 || nicIdx >= (int)closestGpusForNic.size()) continue; + if (!closestGpusForNic[nicIdx].empty()) closestGpusForNic[nicIdx] += ","; + closestGpusForNic[nicIdx] += std::to_string(j); } + } + for (int i = 0; i < ibvDeviceList.size(); i++) { + std::string closestGpusStr = closestGpusForNic[i]; printf(" %-3d | %-11s | %-6s | %-12s | %-4d | %-14s | %-9s | %-20s\n", i, ibvDeviceList[i].name.c_str(), ibvDeviceList[i].hasActivePort ? "Yes" : "No", @@ -89,7 +83,13 @@ void DisplaySingleRankTopology(bool outputToCsv) } else { printf("\nDetected Topology:\n"); printf("==================\n"); - printf(" %d configured CPU NUMA node(s) [%d total]\n", numCpus, numa_max_node() + 1); + int hiddenNuma = numa_num_configured_nodes() - numCpus; + if (hiddenNuma > 0) { + printf(" %d configured CPU NUMA node(s) [%d total] (%d core-less node(s) hidden; set TB_SHOW_ALL_NUMA=1 to show)\n", + numCpus, numa_max_node() + 1, hiddenNuma); + } else { + printf(" %d configured CPU NUMA node(s) [%d total]\n", numCpus, numa_max_node() + 1); + } printf(" %d GPU device(s)\n", numGpus); printf(" %d Supported NIC device(s)\n", numNics); } @@ -121,8 +121,9 @@ void DisplaySingleRankTopology(bool outputToCsv) if (numa_node_of_cpu(j) == nodeI) numCpuCores++; printf(" %5d %c", numCpuCores, sep); + // GetClosestCpuNumaToGpu returns a logical CPU index, so compare against i (logical) for (int j = 0; j < numGpus; j++) { - if (TransferBench::GetClosestCpuNumaToGpu(j) == nodeI) { + if (TransferBench::GetClosestCpuNumaToGpu(j) == i) { printf(" %d", j); } } @@ -191,13 +192,23 @@ void DisplaySingleRankTopology(bool outputToCsv) char pciBusId[20]; HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, i)); - printf(" %-11s %c %-4d %c %-4d %c %-4d %c %-4d %c %-4d\n", + // Space-separated so the field stays a single column when sep is a comma (CSV mode), + // matching how the "Closest GPU(s)" column above is emitted + std::vector nicsForGpu; + TransferBench::GetClosestNicsToGpu(nicsForGpu, i); + std::string nicStr; + for (int nicIdx : nicsForGpu) { + if (!nicStr.empty()) nicStr += ' '; + nicStr += std::to_string(nicIdx); + } + if (nicStr.empty()) nicStr = "-1"; + printf(" %-11s %c %-4d %c %-4d %c %-4d %c %-4d %c %s\n", pciBusId, sep, TransferBench::GetNumSubExecutors({EXE_GPU_GFX, i}), sep, TransferBench::GetClosestCpuNumaToGpu(i), sep, TransferBench::GetNumExecutorSubIndices({EXE_GPU_DMA, i}), sep, TransferBench::GetNumExecutorSubIndices({EXE_GPU_GFX, i}), sep, - TransferBench::GetClosestNicToGpu(i)); + nicStr.c_str()); } } #endif diff --git a/src/client/Utilities.hpp b/src/client/Utilities.hpp index dafa7541..ad3a5641 100644 --- a/src/client/Utilities.hpp +++ b/src/client/Utilities.hpp @@ -128,6 +128,9 @@ namespace TransferBench::Utils // Return the number of homogenous groups of ranks int GetNumRankGroups(); + // Return true if all ranks have the same number of GPUs + bool AllRanksHaveSameGpuCount(); + // Helper function for pod membership RankPerPodMap& GetRankPerPodMap(); @@ -414,6 +417,16 @@ namespace TransferBench::Utils return GetRankGroupMap().size(); } + bool AllRanksHaveSameGpuCount() + { + int const numRanks = TransferBench::GetNumRanks(); + if (numRanks <= 1) return true; + int const gpuCount = TransferBench::GetNumExecutors(EXE_GPU_GFX, 0); + for (int rank = 1; rank < numRanks; rank++) + if (TransferBench::GetNumExecutors(EXE_GPU_GFX, rank) != gpuCount) return false; + return true; + } + RankPerPodMap& GetRankPerPodMap() { static RankPerPodMap pods; @@ -439,6 +452,7 @@ namespace TransferBench::Utils case EXE_NIC: return "NIC"; case EXE_NIC_NEAREST: return "NIC"; case EXE_GPU_BDMA: return "BMA"; + case EXE_GPU_TDM: return "TDM"; default: return "N/A"; } } @@ -512,6 +526,7 @@ namespace TransferBench::Utils va_start(args, format); vprintf(format, args); va_end(args); + fflush(stdout); } } diff --git a/src/header/TransferBench.hpp b/src/header/TransferBench.hpp index 259fc4dc..d2a59f4b 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 @@ -33,10 +37,12 @@ THE SOFTWARE. #include #include #include +#include #include #include #include #include +#include #include // If not found, try installing libnuma-dev (e.g apt-get install libnuma-dev) #include #include @@ -77,7 +83,10 @@ THE SOFTWARE. #ifdef AMD_SMI_ENABLED #include "amd_smi/amdsmi.h" #endif + #endif + +#include "tdmCopy.h" /// @endcond // Batched DMA executor is only supported with HIP >= 7.1 and CUDA 12.8 @@ -92,7 +101,7 @@ namespace TransferBench using std::set; using std::vector; - constexpr char VERSION[] = "1.69"; + constexpr char VERSION[] = "1.70"; /** * Enumeration of supported Executor types @@ -107,10 +116,11 @@ namespace TransferBench EXE_NIC = 3, ///< NIC RDMA executor (subExecutor = queue pair) EXE_NIC_NEAREST = 4, ///< NIC RDMA nearest executor (subExecutor = queue pair) EXE_GPU_BDMA = 5, ///< GPU Batched SDMA executor (subExecutor = batch item) + EXE_GPU_TDM = 6, ///< GPU TDM executor (subExecutor = threadblock/CU) }; - char const ExeTypeStr[7] = "CGDINB"; + char const ExeTypeStr[8] = "CGDINBT"; inline bool IsCpuExeType(ExeType e){ return e == EXE_CPU; } - inline bool IsGpuExeType(ExeType e){ return e == EXE_GPU_GFX || e == EXE_GPU_DMA || e == EXE_GPU_BDMA; } + inline bool IsGpuExeType(ExeType e){ return e == EXE_GPU_GFX || e == EXE_GPU_DMA || e == EXE_GPU_BDMA || e == EXE_GPU_TDM; } inline bool IsNicExeType(ExeType e){ return e == EXE_NIC || e == EXE_NIC_NEAREST; } /** @@ -209,7 +219,9 @@ namespace TransferBench int numSubIterations = 1; ///< Number of sub-iterations per iteration int numWarmups = 3; ///< Number of un-timed warmup iterations to perform int recordPerIteration = 0; ///< Record per-iteration timing information + int useHipEvents = 1; ///< Use HIP events for timing Executors that support it int useInteractive = 0; ///< Pause for user-input before starting transfer loop + int useMultiStream = 0; ///< Split GFX/TDM Transfers into separate kernel launches in separate stream }; /** @@ -223,6 +235,7 @@ namespace TransferBench vector fillPattern = {}; ///< Pattern of floats used to fill source data vector fillCompress = {}; ///< Customized data patterns (overrides fillPattern if non-empty) int validateDirect = 0; ///< Validate GPU results directly instead of copying to host + int validateOnDevice = 0; ///< Validate GPU dst/src memory via an on-device kernel instead of host memcmp int validateSource = 0; ///< Validate src GPU memory immediately after preparation }; @@ -239,8 +252,6 @@ namespace TransferBench int seType = 0; ///< SubExecutor granularity type (0=threadblock, 1=warp) int temporalMode = 0; ///< Non-temporal load/store mode 0=none, 1=load, 2=store, 3=both int unrollFactor = 4; ///< GFX-kernel unroll factor - int useHipEvents = 1; ///< Use HIP events for timing GFX Executor - int useMultiStream = 0; ///< Use multiple streams for GFX int useSingleTeam = 0; ///< Team all subExecutors across the data array int waveOrder = 0; ///< GFX-kernel wavefront ordering int wordSize = 4; ///< GFX-kernel packed data size (4=dwordx4, 2=dwordx2, 1=dwordx1) @@ -251,7 +262,6 @@ namespace TransferBench */ struct DmaOptions { - int useHipEvents = 1; ///< Use HIP events for timing DMA Executor int useHsaCopy = 0; ///< Use HSA copy instead of HIP copy to perform DMA }; @@ -275,6 +285,12 @@ namespace TransferBench int useNuma = 0; ///< Switch to closest numa thread for execution }; + struct TdmOptions + { + int blockOrder = 0; ///< Determines how threadblocks are ordered (0=sequential, 1=interleaved, 2=random) + int blockSize = 256; ///< Size of each threadblock + int ldsBytes = 0; ///< Amount of __shared__ memory per threadblock to use as bounce buffer (0 = device max) + }; /** * 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 }; /** @@ -474,6 +491,14 @@ namespace TransferBench */ int GetClosestCpuNumaToGpu(int gpuIndex, int targetRank = -1); + /** + * Returns the physical NUMA node id backing a logical CPU NUMA index + * + * @param[in] cpuIndex Logical CPU NUMA index (as exposed by GetNumExecutors(EXE_CPU)) + * @returns Physical NUMA node id, or cpuIndex unchanged if out of range + */ + int GetCpuNumaPhysicalNode(int cpuIndex); + /** * Returns the index of the NUMA node closest to the given NIC * @@ -620,6 +645,7 @@ namespace TransferBench // Enumerations #define hipDeviceAttributeClockRate cudaDevAttrClockRate #define hipDeviceAttributeMultiprocessorCount cudaDevAttrMultiProcessorCount + #define hipDeviceAttributeMaxSharedMemoryPerBlock cudaDevAttrMaxSharedMemoryPerBlock #define hipDeviceAttributeWarpSize cudaDevAttrWarpSize #define hipErrorPeerAccessAlreadyEnabled cudaErrorPeerAccessAlreadyEnabled #define hipFuncCachePreferShared cudaFuncCachePreferShared @@ -649,6 +675,7 @@ namespace TransferBench #define hipGetDeviceCount cudaGetDeviceCount #define hipGetDeviceProperties cudaGetDeviceProperties #define hipGetErrorString cudaGetErrorString + #define hipGetLastError cudaGetLastError #define hipHostFree cudaFreeHost #define hipHostMalloc cudaMallocHost #define hipMalloc cudaMalloc @@ -700,30 +727,75 @@ namespace TransferBench // Helper macro functions //========================================================================================== -// Macro for collecting CU/SM GFX kernel is running on -#if defined(__GFX9__) - #define GetHwId(hwId) asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s" (hwId)) +// Returns xccId and a unified cuId for the current wavefront. +// cuId encoding is arch-specific (see comments) but is always a dense index +// suitable for CU-set tracking. +__device__ __forceinline__ void GetXccHwId(uint32_t& xccId, uint32_t& cuId) +{ +#if defined(__gfx942__) || defined(__gfx950__) + // CDNA3: HW_REG_HW_ID (code 4) + HW_REG_XCC_ID (code 20) + // CDNA3 ISA §5.8 Table 19, §3.12 Table 6 + uint32_t hwId = 0, xccReg = 0; + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s"(hwId)); + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s"(xccReg)); + xccId = xccReg & 0xF; + cuId = (((hwId >> 12) & 1) << 5) // SH_ID + | (((hwId >> 8) & 15) << 2) // CU_ID + | ((hwId >> 13) & 3); // SE_ID + +#elif defined(__gfx1250__) + // CDNA5: HW_ID1 (code 23) + RTN_GET_SE_HW_ID (0x87) + // CDNA5 ISA §3.4.9, §5.4 Table 19 + uint32_t hwId = 0, seHwId = 0; + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s"(hwId)); + asm volatile("s_sendmsg_rtn_b32 %0, 0x87\ns_wait_kmcnt 0" : "=s"(seHwId)); + xccId = (seHwId >> 16) & 0xF; // Virtual_XCC_ID [19:16] + cuId = ((seHwId & 0xF) << 4) // SE_ID [3:0] (gfx1250: 2 SEs → 1 bit) + | (((hwId >> 16) & 0x1) << 3) // SA_ID [16] (gfx1250: 2 SAs → 1 bit) + | ((hwId >> 10) & 0x7); // WGP_ID [12:10] (gfx1250: 8 WGPs per SA → 3 bits) + +#elif defined(__GFX9__) + // Other GFX9 (gfx90a/MI200, gfx908, gfx906) — single die, no XCC + // CDNA2 ISA §3.12 Table 6 + uint32_t hwId = 0; + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s"(hwId)); + xccId = 0; + cuId = (((hwId >> 12) & 1) << 5) + | (((hwId >> 8) & 15) << 2) + | ((hwId >> 13) & 3); + #elif defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__) - #define GetHwId(hwId) asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s" (hwId)) + // RDNA2/3/4 (non-CDNA5) — HW_ID1 present, no XCC + uint32_t hwId = 0; + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s"(hwId)); + xccId = 0; + cuId = hwId; + #elif defined(__NVCC__) - #define GetHwId(hwId) asm("mov.u32 %0, %smid;" : "=r"(hwId)) + xccId = 0; + asm("mov.u32 %0, %smid;" : "=r"(cuId)); #else - #define GetHwId(hwId) hwId = 0 + xccId = 0; + cuId = 0; #endif +} -// Macro for collecting XCC GFX kernel is running on +// Returns only the XCC ID for the current wavefront; cheaper than GetXccHwId +// when the cuId is not needed. +__device__ __forceinline__ uint32_t GetXccId() +{ #if defined(__gfx942__) || defined(__gfx950__) -#define GetXccId(val) asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s" (val)) -#elif defined(__GFX12__) -#define GetXccId(val) \ - { asm volatile ("s_sendmsg_rtn_b32 %0, 0x87 \n" \ - "s_wait_kmcnt 0" \ - : "=s" (val)); \ - val = ((val >> 16) & 0xF); \ - } + uint32_t xccReg = 0; + asm volatile("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s"(xccReg)); + return xccReg & 0xF; +#elif defined(__gfx1250__) + uint32_t seHwId = 0; + asm volatile("s_sendmsg_rtn_b32 %0, 0x87\ns_wait_kmcnt 0" : "=s"(seHwId)); + return (seHwId >> 16) & 0xF; // Virtual_XCC_ID [19:16] #else -#define GetXccId(val) val = 0 + return 0; #endif +} // Error check macro (NOTE: This will return even for ERR_WARN) #define ERR_CHECK(cmd) \ @@ -915,7 +987,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) /** * Barrier that all ranks must arrive at before proceeding */ - void Barrier(); + void Barrier() const; /** * Send data to a single destination rank @@ -1039,6 +1111,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::string GetExecutorName(ExeDevice exeDevice) const; int NicIsActive(int nicIndex, int targetRank) const; + // Translate a logical CPU NUMA index (as exposed to users) into the physical + // NUMA node id used by libnuma / HSA. Returns the index unchanged if out of range. + int GetCpuPhysicalNode(int logicalIdx) const; + // Translate a physical NUMA node id into its logical CPU index, or -1 if the node + // is not exposed (e.g. a core-less node skipped by default). + int GetCpuLogicalNode(int physicalNode) const; + #if !defined(__NVCC__) ErrResult GetHsaAgent(ExeDevice const& exeDevice, hsa_agent_t& agent) const; ErrResult GetHsaAgent(MemDevice const& memDevice, hsa_agent_t& agent) const; @@ -1070,6 +1149,12 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::vector gpuAgents; #endif + // CPU NUMA remapping (logical index <-> physical NUMA node) + // By default core-less NUMA nodes are skipped; TB_SHOW_ALL_NUMA=1 exposes every node. + bool showAllNuma = false; ///< Expose all configured NUMA nodes (legacy behavior) + std::vector cpuNumaMap; ///< logical index -> physical NUMA node + std::map cpuNumaRevMap; ///< physical NUMA node -> logical index + int commMode; ///< Communication mode #ifdef MPI_COMM_ENABLED @@ -1080,8 +1165,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Socket related std::string masterAddr; ///< Rank 0 master address int masterPort; ///< Rank 0 master port - std::vector sockets; ///< Master list of sockets - int listenSocket; ///< Master listener socket + std::vector sockets; ///< Master list of sockets + mutable std::vector socketEpoch; ///< Per-socket send/recv epoch counter + int listenSocket; ///< Master listener socket + int sendUsleep = 0; ///< Microseconds to sleep after each SendData (TB_SEND_USLEEP) // Topology related struct RankTopology @@ -1106,6 +1193,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) void SetupSocketCommunicator(); void SetupMpiCommunicator(); void CollectPodMembership(char* ppodId, int64_t& vpodId); + void BuildCpuNumaMap(); + // Return the logical index of the exposed CPU NUMA node nearest (by numa_distance) to a + // physical node; used when the physical node itself is not exposed (e.g. core-less). + int GetClosestLogicalCpu(int physicalNode) const; void GetRankTopology(RankTopology& topo); void CollectTopology(); std::string GetCpuName() const; @@ -1530,9 +1621,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) deviceIdx = GetClosestCpuNumaToGpu(memDevice.memIndex); } + // For CPU memory, deviceIdx is a logical CPU NUMA index; translate to the physical + // NUMA node for all libnuma operations (policy, allocation, page verification). + int cpuPhysNode = IsCpuMemType(memType) ? System::Get().GetCpuPhysicalNode(deviceIdx) : deviceIdx; + if (IsCpuMemType(memType)) { // Set NUMA policy prior to call to hipHostMalloc - numa_set_preferred(deviceIdx); + numa_set_preferred(cpuPhysNode); } else if (IsGpuMemType(memType)) { // Switch to the appropriate GPU // IMP: if the remapping above changes, remember to modify this! @@ -1574,7 +1669,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (IsCpuMemType(memType)) { memset(*memPtr, 0, roundedUpBytes); // Check that the allocated pages are actually on the correct NUMA node - ERR_CHECK(CheckPages((char*)*memPtr, roundedUpBytes, deviceIdx)); + ERR_CHECK(CheckPages((char*)*memPtr, roundedUpBytes, cpuPhysNode)); numa_set_preferred(-1); } else if (IsGpuMemType(memType)) { ERR_CHECK(hipMemset(*memPtr, 0, numBytes)); @@ -1621,12 +1716,12 @@ const auto& AmdSmiFabricInfoV1(const T& info) #endif #endif } else if (memType == MEM_CPU_UNPINNED) { - *memPtr = numa_alloc_onnode(numBytes, deviceIdx); + *memPtr = numa_alloc_onnode(numBytes, cpuPhysNode); } // Check that the allocated pages are actually on the correct NUMA node memset(*memPtr, 0, numBytes); - ERR_CHECK(CheckPages((char*)*memPtr, numBytes, deviceIdx)); + ERR_CHECK(CheckPages((char*)*memPtr, numBytes, cpuPhysNode)); // Reset to default numa mem policy numa_set_preferred(-1); @@ -1917,7 +2012,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) return {ERR_FATAL, "CPU index must be between 0 and %d (instead of %d) on rank %d", numCpus - 1, memDevice.memIndex, memDevice.memRank}; - if (GetRank() == memDevice.memRank && !numa_bitmask_isbitset(numa_get_mems_allowed(), memDevice.memIndex)) { + if (GetRank() == memDevice.memRank && + !numa_bitmask_isbitset(numa_get_mems_allowed(), + System::Get().GetCpuPhysicalNode(memDevice.memIndex))) { return {ERR_FATAL, "CPU %d on rank %d cannot allocate memory due to process memory policy/cpuset", memDevice.memIndex, memDevice.memRank}; } @@ -1934,7 +2031,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (actualNumaIdx == -1) { return {ERR_FATAL, "Unable to determine closest NUMA node for GPU %d on rank %d", memDevice.memIndex, memDevice.memRank}; } - if (GetRank() == memDevice.memRank && !numa_bitmask_isbitset(numa_get_mems_allowed(), actualNumaIdx)) + if (GetRank() == memDevice.memRank && + !numa_bitmask_isbitset(numa_get_mems_allowed(), + System::Get().GetCpuPhysicalNode(actualNumaIdx))) return {ERR_FATAL, "CPU %d on rank %d cannot allocate memory due to process memory policy/cpuset", memDevice.memIndex, memDevice.memRank}; } @@ -1964,7 +2063,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (general.numSubIterations != cfg.general.numSubIterations) ADD_ERROR("cfg.general.numSubIterations"); if (general.numWarmups != cfg.general.numWarmups) ADD_ERROR("cfg.general.numWarmups"); if (general.recordPerIteration != cfg.general.recordPerIteration) ADD_ERROR("cfg.general.recordPerIteration"); + if (general.useHipEvents != cfg.general.useHipEvents) ADD_ERROR("cfg.general.useHipEvents"); if (general.useInteractive != cfg.general.useInteractive) ADD_ERROR("cfg.general.useInteractive"); + if (general.useMultiStream != cfg.general.useMultiStream) ADD_ERROR("cfg.general.useMultiStream"); } // Compare data options @@ -2030,8 +2131,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (gfx.seType != cfg.gfx.seType) ADD_ERROR("cfg.gfx.seType"); if (gfx.temporalMode != cfg.gfx.temporalMode) ADD_ERROR("cfg.gfx.temporalMode"); if (gfx.unrollFactor != cfg.gfx.unrollFactor) ADD_ERROR("cfg.gfx.unrollFactor)"); - if (gfx.useHipEvents != cfg.gfx.useHipEvents) ADD_ERROR("cfg.gfx.useHipEvents"); - if (gfx.useMultiStream != cfg.gfx.useMultiStream) ADD_ERROR("cfg.gfx.useMultiStream"); if (gfx.useSingleTeam != cfg.gfx.useSingleTeam) ADD_ERROR("cfg.gfx.useSingleTeam"); if (gfx.waveOrder != cfg.gfx.waveOrder) ADD_ERROR("cfg.gfx.waveOrder"); if (gfx.wordSize != cfg.gfx.wordSize) ADD_ERROR("cfg.gfx.wordSize"); @@ -2041,7 +2140,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) { DmaOptions dma = cfg.dma; System::Get().Broadcast(root, sizeof(dma), &dma); - if (dma.useHipEvents != cfg.dma.useHipEvents) ADD_ERROR("cfg.dma.useHipEvents"); if (dma.useHsaCopy != cfg.dma.useHsaCopy) ADD_ERROR("cfg.dma.useHsaCopy"); } @@ -2064,6 +2162,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (nic.useNuma != cfg.nic.useNuma) ADD_ERROR("cfg.nic.useNuma"); } + // Compare TDM options + { + TdmOptions tdm = cfg.tdm; + System::Get().Broadcast(root, sizeof(tdm), &tdm); + if (tdm.blockOrder != cfg.tdm.blockOrder) ADD_ERROR("cfg.tdm.blockOrder"); + if (tdm.blockSize != cfg.tdm.blockSize) ADD_ERROR("cfg.tdm.blockSize"); + if (tdm.ldsBytes != cfg.tdm.ldsBytes) ADD_ERROR("cfg.tdm.ldsBytes"); + } #undef ADD_ERROR } @@ -2105,7 +2211,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) errors.push_back({ERR_FATAL, "[gfx.blockOrder] must be 0 for sequential, 1 for interleaved, or 2 for random"}); - if (cfg.gfx.useMultiStream && cfg.gfx.blockOrder > 0) + if (cfg.general.useMultiStream && cfg.gfx.blockOrder > 0) errors.push_back({ERR_WARN, "[gfx.blockOrder] will be ignored when running in multi-stream mode"}); if (cfg.gfx.blockSize < 0 || cfg.gfx.blockSize % 64 || cfg.gfx.blockSize > MAX_BLOCKSIZE) @@ -2163,6 +2269,34 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } + // Check TDM options (TDM-capable hardware: gfx1250 on AMD, sm_90+ on NVIDIA) + 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.ldsBytes < 0) + errors.push_back({ERR_FATAL, "[tdm.ldsBytes] must be positive or 0"}); + else { + int const numGpus = GetNumExecutors(EXE_GPU_TDM); + for (int i = 0; i < numGpus; i++) { + int deviceMax = 0; + if (hipDeviceGetAttribute(&deviceMax, hipDeviceAttributeMaxSharedMemoryPerBlock, i) == hipSuccess) { + if (cfg.tdm.ldsBytes > deviceMax) { + errors.push_back({ERR_FATAL, + "[tdm.ldsBytes] (%d) exceeds device max shared memory per block (%d)", + cfg.tdm.ldsBytes, deviceMax}); + } else if (deviceMax == 0) { + errors.push_back({ERR_FATAL, + "TDM Executor requires shared memory but GPU %d reports none available\n", i}); + } + } else { + errors.push_back({ERR_FATAL, + "Unable to query max amount of shared memory per block on GPU %d\n", i}); + } + } + } + // Check NIC options if (IsIbvSymbolsReady()) { if (cfg.nic.chunkBytes == 0 || (cfg.nic.chunkBytes % 4 != 0)) { @@ -2296,7 +2430,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) // 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_DMA) { // GPU DMA-Executor ignores subexecutors 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, @@ -2354,6 +2489,33 @@ const auto& AmdSmiFabricInfoV1(const T& info) hasFatalError = true; } break; + case EXE_GPU_TDM: + if (t.srcs.size() != 1 || t.dsts.size() != 1) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel currently requires exactly 1 SRC and 1 DST", i}); + hasFatalError = true; + break; + } + if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel: device index must be between 0 and %d (instead of %d) for rank %d", + i, numExecutors - 1, t.exeDevice.exeIndex, t.exeDevice.exeRank}); + hasFatalError = true; + break; + } + if (t.exeSubIndex != -1) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM executor does not support subindices", i}); + hasFatalError = true; + break; + } + if (!tdm::IsTdmCopySupported(t.exeDevice.exeIndex)) { + errors.push_back({ERR_FATAL, + "Transfer %d: GPU TDM kernel requires TDM-capable hardware (gfx1250 or NVIDIA sm_90+), but GPU %d is not supported", + i, t.exeDevice.exeIndex}); + hasFatalError = true; + } + break; case EXE_GPU_GFX: if (t.exeDevice.exeIndex < 0 || t.exeDevice.exeIndex >= numExecutors) { errors.push_back({ERR_FATAL, @@ -2374,7 +2536,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) 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 } @@ -2717,13 +2878,29 @@ const auto& AmdSmiFabricInfoV1(const T& info) break; } - if (cfg.gfx.useMultiStream && transferCount[exeDevice] > gpuMaxHwQueues) { + if (cfg.general.useMultiStream && transferCount[exeDevice] > gpuMaxHwQueues) { errors.push_back({ERR_WARN, "GPU %d attempting %d parallel transfers, however GPU_MAX_HW_QUEUES only set to %d", exeDevice.exeIndex, transferCount[exeDevice], gpuMaxHwQueues}); } break; } + case EXE_GPU_TDM: + { + int numGpuSubExec = GetNumSubExecutors(exeDevice); + if (totalSubExecs[exeDevice] > numGpuSubExec) + errors.push_back({ERR_WARN, + "TDM %d requests %d total subexecutors however only %d available. " + "Serialization will occur", + exeDevice.exeIndex, totalSubExecs[exeDevice], numGpuSubExec}); + + if (cfg.general.useMultiStream && transferCount[exeDevice] > gpuMaxHwQueues) { + errors.push_back({ERR_WARN, + "TDM %d 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 @@ -2811,6 +2988,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) // For GFX executor SubExecParam* subExecParamGpuPtr; + // For on-device validation (VALIDATE_ON_DEVICE) + vector dstExpectedMem; ///< Per-dst device copy of expected values + vector dstExpectedBytes; ///< Allocated size of each dst expected buffer + vectordstValidateScratch; ///< Per-dst device scratch [0]=count, [1]=firstOffset + vector srcExpectedMem; ///< Per-src device copy of expected values (validateSource) + vector srcExpectedBytes; ///< Allocated size of each src expected buffer + vectorsrcValidateScratch; ///< Per-src device scratch [0]=count, [1]=firstOffset + // For targeted-SDMA #if !defined(__NVCC__) vector dstAgent; ///< DMA destination memory agents @@ -2858,33 +3043,139 @@ const auto& AmdSmiFabricInfoV1(const T& info) vector>> perIterCUs; ///< GFX-Executor only. XCC:CU used per iteration }; + // Persistent pool of worker threads, created once and reused across all iterations to + // avoid the per-iteration thread creation overhead of std::async/std::thread. + // Each worker runs an optional init hook once at startup (used to bind NUMA affinity and + // set the HIP device for the executor the pool serves). + class ThreadPool + { + public: + // numThreads worker threads are spawned immediately. perThreadInit (if provided) runs + // once on each worker before it starts servicing tasks. + ThreadPool(int numThreads, std::function perThreadInit = {}) + { + numThreads = std::max(1, numThreads); + workers.reserve(numThreads); + for (int i = 0; i < numThreads; ++i) + workers.emplace_back(&ThreadPool::WorkerLoop, this, perThreadInit); + } + + ~ThreadPool() + { + { + std::unique_lock lock(mutex); + stop = true; + } + cv.notify_all(); + for (auto& worker : workers) + if (worker.joinable()) worker.join(); + } + + ThreadPool(ThreadPool const&) = delete; + ThreadPool& operator=(ThreadPool const&) = delete; + + // Run fn(i) for i in [0, count), spreading the calls across the worker threads, and block + // until every call has completed. Safe to call repeatedly; no allocation per call. + void ParallelFor(int count, std::function const& fn) + { + if (count <= 0) return; + { + // Publish the new batch. Every worker is woken and runs until the shared index is + // exhausted; completion is signalled once all workers have parked again (not when the + // last task finishes) so a lagging worker can never observe the next batch's reset. + std::unique_lock lock(mutex); + task = &fn; + nextIndex.store(0); + totalTasks = count; + doneWorkers = 0; + ++generation; + } + cv.notify_all(); + + // All work runs on the (NUMA-pinned) worker threads; the caller only waits so that a + // subexecutor's memory traffic is never issued from an unpinned dispatch thread. + std::unique_lock lock(mutex); + doneCv.wait(lock, [this] { return doneWorkers == (int)workers.size(); }); + } + + private: + void WorkerLoop(std::function perThreadInit) + { + if (perThreadInit) perThreadInit(); + + uint64_t lastGeneration = 0; + while (true) { + std::function const* localTask = nullptr; + int localCount = 0; + { + std::unique_lock lock(mutex); + cv.wait(lock, [this, &lastGeneration] { return stop || generation != lastGeneration; }); + if (stop) return; + lastGeneration = generation; + localTask = task; + localCount = totalTasks; + } + + // Claim task indices off the shared counter until exhausted + while (true) { + int idx = nextIndex.fetch_add(1); + if (idx >= localCount) break; + (*localTask)(idx); + } + + // Mark this worker parked; the last one to park wakes the waiting caller + std::unique_lock lock(mutex); + if (++doneWorkers == (int)workers.size()) + doneCv.notify_one(); + } + } + + std::vector workers; + std::mutex mutex; ///< Guards dispatch, generation, completion + std::condition_variable cv; ///< Wakes workers for a new batch + std::condition_variable doneCv; ///< Wakes ParallelFor when batch completes + std::function const* task = nullptr; ///< Current batch task (owned by caller) + std::atomic nextIndex{0}; ///< Next task index to claim + int totalTasks = 0; ///< Size of current batch + int doneWorkers = 0; ///< Workers parked for the current batch + uint64_t generation = 0; ///< Incremented per batch to wake workers + bool stop = false; ///< Set at destruction + }; + // Internal resources allocated per Executor struct ExeInfo { - size_t totalBytes; ///< Total bytes this executor transfers - double totalDurationMsec; ///< Total duration for all iterations for this Executor - int totalSubExecs; ///< Total number of subExecutors to use - bool useSubIndices; ///< Use subexecutor indicies - int numSubIndices; ///< Number of subindices this ExeDevice has - vector subExecParamCpu; ///< Subexecutor parameters for this executor - vector resources; ///< Per-Transfer resources + size_t totalBytes; ///< Total bytes this executor transfers + double totalDurationMsec; ///< Total duration for all iterations for this Executor + int totalSubExecs; ///< Total number of subExecutors to use + bool useSubIndices; ///< Use subexecutor indicies + int numSubIndices; ///< Number of subindices this ExeDevice has + vector subExecParamCpu; ///< Subexecutor parameters for this executor + vector resources; ///< Per-Transfer resources // For GPU-Executors - SubExecParam* subExecParamGpu; ///< GPU copy of subExecutor parameters + SubExecParam* subExecParamGpu; ///< GPU copy of subExecutor parameters bool subExecParamHostAccessible; ///< Host can directly read subExecParamGpu - vector streams; ///< HIP streams to launch on - vector startEvents; ///< HIP start timing event - vector stopEvents; ///< HIP stop timing event - int wallClockRate; ///< (GFX-only) Device wall clock rate - int gfxKernelToUse; ///< (GFX-only) Which GFX kernel to use + vector streams; ///< HIP streams to launch on + vector startEvents; ///< HIP start timing event + vector stopEvents; ///< HIP stop timing event + int wallClockRate; ///< (GFX-only) Device wall clock rate + int gfxKernelToUse; ///< (GFX-only) Which GFX kernel to use + + // For TDM-Executors + uint32_t ldsBytesActual; ///< Actual number of LDS bytes to use as buffer + + // Persistent worker pool servicing this executor's Transfers/subexecutors across all + // iterations. Created in PrepareExecutor (NUMA-pinned), destroyed in TeardownExecutor. + std::unique_ptr pool; }; // Structure to track PCIe topology struct PCIeNode { - std::string address; ///< PCIe address for this PCIe node - std::string description; ///< Description for this PCIe node - std::set children; ///< Children PCIe nodes + std::string address; ///< PCIe address for this PCIe node + std::string description; ///< Description for this PCIe node + std::set children; ///< Children PCIe nodes // Default constructor PCIeNode() : address(""), description("") {} @@ -3111,7 +3402,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) ibvDeviceList.push_back(ibvDevice); } } - ibv_free_device_list(deviceList); + if (deviceList) ibv_free_device_list(deviceList); isInitialized = true; } return ibvDeviceList; @@ -3242,8 +3533,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) return -1; } - // Function to extract the bus number from a PCIe address (domain:bus:device.function) - static int ExtractBusNumber(std::string const& pcieAddress) + // Function to extract the domain number from a PCIe address (domain:bus:device.function) + // All four fields are read (not just the domain) so that an address with too few fields, or + // with a non-hex field, is rejected. The separator characters themselves are consumed but + // not checked, so this is a well-formedness guard rather than strict format validation + static int ExtractDomain(std::string const& pcieAddress) { int domain, bus, device, function; char delimiter; @@ -3256,16 +3550,31 @@ const auto& AmdSmiFabricInfoV1(const T& info) #endif return -1; } - return bus; - } - - // Function to compute the distance between two bus IDs - static int GetBusIdDistance(std::string const& pcieAddress1, - std::string const& pcieAddress2) - { - int bus1 = ExtractBusNumber(pcieAddress1); - int bus2 = ExtractBusNumber(pcieAddress2); - return (bus1 < 0 || bus2 < 0) ? -1 : std::abs(bus1 - bus2); + return domain; + } + + // Computes a proximity distance between two PCIe addresses. Used as a secondary + // tiebreaker when candidates share the same LCA depth in the PCIe tree, and as the + // sole metric in the fallback path when the PCIe tree yields no match at all. + // Returns -1 if either address cannot be parsed. + // + // Same domain (0): devices in one PCIe domain share a root complex, so the LCA tree + // already captures their structural proximity. Bus numbers are firmware-assigned and + // do not reliably track physical closeness, so they are deliberately NOT used to + // discriminate within a domain -- doing so would break ties between NICs that are + // genuinely equidistant from a GPU (e.g. two NICs hanging off the same root complex + // at different bus numbers), which is exactly the case this metric must preserve. + // + // Cross domain (|delta domain|): any non-zero value ranks behind every same-domain + // candidate. The magnitude only provides a deterministic ordering among cross-domain + // candidates; it carries no physical meaning. + static int GetDomainDistance(std::string const& pcieAddress1, + std::string const& pcieAddress2) + { + int domain1 = ExtractDomain(pcieAddress1); + int domain2 = ExtractDomain(pcieAddress2); + if (domain1 < 0 || domain2 < 0) return -1; + return std::abs(domain1 - domain2); } // Given a target busID and a set of candidate devices, returns a set of indices @@ -3285,11 +3594,15 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (!lca) continue; int depth = GetLcaDepth(lca->address, GetPCIeTreeRoot()); - int currDistance = GetBusIdDistance(targetBusId, candidateBusId); + int currDistance = GetDomainDistance(targetBusId, candidateBusId); - // When more than one LCA match is found, choose the one with smallest busId difference - // NOTE: currDistance could be -1, which signals problem with parsing, however still - // remains a valid "closest" candidate, so is included + // A candidate whose address could not be parsed (-1) remains eligible, but treat its + // distance as the largest possible so it can never outrank a candidate whose distance + // is actually known. It can still be selected when nothing else matches at this depth, + // and still ties with other unparseable candidates. + if (currDistance < 0) currDistance = std::numeric_limits::max(); + + // When more than one LCA match is found, choose the one with smallest domain difference if (depth > maxDepth || (depth == maxDepth && depth >= 0 && currDistance < minDistance)) { maxDepth = depth; matches.clear(); @@ -3924,6 +4237,100 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } + // Scans output vs expected byte-by-byte and returns a summary of contiguous mismatch + // ranges (byte indices, inclusive). Reports up to 5 ranges then "..." if more exist, + // followed by the total mismatch byte count. + static std::string ByteMismatchSummary(void const* output, void const* expected, size_t numBytes) + { + unsigned char const* out = static_cast(output); + unsigned char const* exp = static_cast(expected); + + // Collect all contiguous mismatch ranges in one pass + struct Range { size_t start, end; }; + std::vector ranges; + size_t totalMismatch = 0; + size_t i = 0; + while (i < numBytes) { + if (out[i] != exp[i]) { + size_t start = i; + while (i < numBytes && out[i] != exp[i]) ++i; + totalMismatch += i - start; + ranges.push_back({start, i - 1}); + } else { + ++i; + } + } + + std::string result; + int const maxReport = 5; + for (int r = 0; r < (int)ranges.size() && r < maxReport; ++r) { + if (r > 0) result += ", "; + char buf[64]; + snprintf(buf, sizeof(buf), "%zu-%zu", ranges[r].start, ranges[r].end); + result += buf; + } + if ((int)ranges.size() > maxReport) result += ", ..."; + char summary[64]; + snprintf(summary, sizeof(summary), " (total %zu mismatched bytes)", totalMismatch); + result += summary; + return result; + } + + // On-device validation kernel: byte-compares a buffer against expected values, recording the + // number of mismatched bytes and the offset of the first mismatch. Compiles under HIP and CUDA. + __global__ void GpuValidateKernel(unsigned char const* __restrict__ data, + unsigned char const* __restrict__ expected, + size_t numBytes, + unsigned long long* result) // [0]=count, [1]=firstOffset + { + size_t const stride = (size_t)gridDim.x * blockDim.x; + for (size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; i < numBytes; i += stride) { + if (data[i] != expected[i]) { + atomicAdd(result, 1ULL); + atomicMin(result + 1, (unsigned long long)i); + } + } + } + + // Launches GpuValidateKernel on the given device and returns the mismatch count / first offset. + // scratch is a 2-element device buffer ([count, firstOffset]) allocated on the same device. + static ErrResult ValidateBufferOnDevice(int deviceIndex, + float const* devData, + float const* devExpected, + size_t numBytes, + unsigned long long* scratch, + unsigned long long* outCount, + unsigned long long* outFirst) + { + ERR_CHECK(hipSetDevice(deviceIndex)); + + // Reset scratch: count=0, firstOffset=numBytes (sentinel meaning "no mismatch") + unsigned long long init[2] = {0ULL, (unsigned long long)numBytes}; + ERR_CHECK(hipMemcpy(scratch, init, sizeof(init), hipMemcpyHostToDevice)); + + int const blockSize = 256; + size_t const numBlocks = (numBytes + blockSize - 1) / blockSize; + int const gridSize = (int)std::min(numBlocks, 4096); + dim3 const grid(gridSize > 0 ? gridSize : 1); + dim3 const block(blockSize); + +#if defined(__NVCC__) + GpuValidateKernel<<>>((unsigned char const*)devData, + (unsigned char const*)devExpected, numBytes, scratch); +#else + hipLaunchKernelGGL(GpuValidateKernel, grid, block, 0, 0, + (unsigned char const*)devData, (unsigned char const*)devExpected, + numBytes, scratch); +#endif + ERR_CHECK(hipDeviceSynchronize()); + + unsigned long long out[2]; + ERR_CHECK(hipMemcpy(out, scratch, sizeof(out), hipMemcpyDeviceToHost)); + *outCount = out[0]; + *outFirst = out[0] ? out[1] : 0ULL; + return ERR_NONE; + } + // Checks that destination buffers match expected values static ErrResult ValidateAllTransfers(ConfigOptions const& cfg, vector const& transfers, @@ -3939,48 +4346,59 @@ const auto& AmdSmiFabricInfoV1(const T& info) for (auto rss : transferResources) { int transferIdx = rss->transferIdx; Transfer const& t = transfers[transferIdx]; - size_t N = t.numBytes / sizeof(float); float const* expected = dstReference[t.srcs.size()].data(); for (int dstIdx = 0; dstIdx < (int)rss->dstMem.size(); dstIdx++) { // Validation is only done on the rank the destination memory is on if (t.dsts[dstIdx].memRank != GetRank()) continue; - if (IsCpuMemType(t.dsts[dstIdx].memType) || cfg.data.validateDirect) { - output = (rss->dstMem[dstIdx]) + initOffset; + + ErrResult dstErr = ERR_NONE; + if (cfg.data.validateOnDevice && IsGpuMemType(t.dsts[dstIdx].memType)) { + // Compare on the GPU against the pre-uploaded expected buffer; only a tiny result copies back if (verbose) { - System::Get().Log("[INFO] Validation: transfer %d DST[%d] direct read from %s%d %zu bytes ptr=%p\n", + System::Get().Log("[INFO] Validation on-device: transfer %d DST[%d] %s%d %zu bytes ptr=%p\n", transferIdx, dstIdx, GetMemTypeName(t.dsts[dstIdx].memType), t.dsts[dstIdx].memIndex, - t.numBytes, output); + t.numBytes, (rss->dstMem[dstIdx]) + initOffset); + } + unsigned long long count = 0, firstOff = 0; + ERR_CHECK(ValidateBufferOnDevice(t.dsts[dstIdx].memIndex, + (rss->dstMem[dstIdx]) + initOffset, + rss->dstExpectedMem[dstIdx], t.numBytes, + rss->dstValidateScratch[dstIdx], &count, &firstOff)); + if (count) { + dstErr = {ERR_FATAL, "Transfer %d: Mismatch at destination %d on rank %d: %llu mismatched bytes, first at offset %llu", + transferIdx, dstIdx, t.dsts[dstIdx].memRank, count, firstOff}; } } else { - ERR_CHECK(hipSetDevice(t.dsts[dstIdx].memIndex)); - if (verbose) { - System::Get().Log("[INFO] Validation memcpy: transfer %d DST[%d] %s%d->host %zu bytes src=%p dst=%p\n", - transferIdx, dstIdx, - GetMemTypeName(t.dsts[dstIdx].memType), t.dsts[dstIdx].memIndex, - t.numBytes, - (rss->dstMem[dstIdx]) + initOffset, - outputBuffer.data()); + if (IsCpuMemType(t.dsts[dstIdx].memType) || cfg.data.validateDirect) { + output = (rss->dstMem[dstIdx]) + initOffset; + if (verbose) { + System::Get().Log("[INFO] Validation: transfer %d DST[%d] direct read from %s%d %zu bytes ptr=%p\n", + transferIdx, dstIdx, + GetMemTypeName(t.dsts[dstIdx].memType), t.dsts[dstIdx].memIndex, + t.numBytes, output); + } + } else { + ERR_CHECK(hipSetDevice(t.dsts[dstIdx].memIndex)); + if (verbose) { + System::Get().Log("[INFO] Validation memcpy: transfer %d DST[%d] %s%d->host %zu bytes src=%p dst=%p\n", + transferIdx, dstIdx, + GetMemTypeName(t.dsts[dstIdx].memType), t.dsts[dstIdx].memIndex, + t.numBytes, + (rss->dstMem[dstIdx]) + initOffset, + outputBuffer.data()); + } + ERR_CHECK(hipMemcpy(outputBuffer.data(), (rss->dstMem[dstIdx]) + initOffset, t.numBytes, hipMemcpyDefault)); + ERR_CHECK(hipDeviceSynchronize()); + output = outputBuffer.data(); } - ERR_CHECK(hipMemcpy(outputBuffer.data(), (rss->dstMem[dstIdx]) + initOffset, t.numBytes, hipMemcpyDefault)); - ERR_CHECK(hipDeviceSynchronize()); - output = outputBuffer.data(); - } - ErrResult dstErr = ERR_NONE; - if (memcmp(output, expected, t.numBytes)) { - // Difference found - find first error - for (size_t i = 0; i < N; i++) { - if (output[i] != expected[i]) { - dstErr = {ERR_FATAL, "Transfer %d: Unexpected mismatch at index %lu of destination %d on rank %d: Expected %10.5f Actual: %10.5f", - transferIdx, i, dstIdx, t.dsts[dstIdx].memRank, expected[i], output[i]}; - break; - } + if (memcmp(output, expected, t.numBytes)) { + std::string ranges = ByteMismatchSummary(output, expected, t.numBytes); + dstErr = {ERR_FATAL, "Transfer %d: Mismatch at destination %d on rank %d: bytes %s", + transferIdx, dstIdx, t.dsts[dstIdx].memRank, ranges.c_str()}; } - if (dstErr.errType == ERR_NONE) - // memcmp found a difference but float != didn't (e.g. +0.0f vs -0.0f bit pattern) - dstErr = {ERR_FATAL, "Transfer %d: Unexpected output mismatch for destination %d", transferIdx, dstIdx}; } if (verbose) @@ -4358,13 +4776,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Prepare additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) - && exeDevice.exeRank == localRank) { + if (IsGpuExeType(exeDevice.exeType) && exeDevice.exeRank == localRank) { ERR_CHECK(hipSetDevice(exeDevice.exeIndex)); // Determine how many streams to use int const numStreamsToUse = (exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA || - (exeDevice.exeType == EXE_GPU_GFX && cfg.gfx.useMultiStream)) + (cfg.general.useMultiStream && (exeDevice.exeType == EXE_GPU_GFX || + exeDevice.exeType == EXE_GPU_TDM))) ? exeInfo.resources.size() : 1; exeInfo.streams.resize(numStreamsToUse); @@ -4382,7 +4800,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { + if (cfg.general.useHipEvents) { exeInfo.startEvents.resize(numStreamsToUse); exeInfo.stopEvents.resize(numStreamsToUse); for (int i = 0; i < numStreamsToUse; ++i) { @@ -4390,10 +4808,23 @@ const auto& AmdSmiFabricInfoV1(const T& info) ERR_CHECK(hipEventCreate(&exeInfo.stopEvents[i])); } } + + // Determine how much shared memory to use for TDM + if (exeDevice.exeType == EXE_GPU_TDM) { + if (cfg.tdm.ldsBytes == 0) { + int ldsMaxBytes; + ERR_CHECK(hipDeviceGetAttribute(&ldsMaxBytes, + hipDeviceAttributeMaxSharedMemoryPerBlock, exeDevice.exeIndex)); + exeInfo.ldsBytesActual = static_cast(ldsMaxBytes); + } else { + exeInfo.ldsBytesActual = cfg.tdm.ldsBytes; + } + } } - // 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__) @@ -4415,7 +4846,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) exeDevice.exeIndex)); #endif int transferOffset = 0; - if (cfg.gfx.useMultiStream || cfg.gfx.blockOrder == 0) { + if (cfg.general.useMultiStream || cfg.gfx.blockOrder == 0) { // Threadblocks are ordered sequentially one transfer at a time for (auto& rss : exeInfo.resources) { rss.subExecParamGpuPtr = exeInfo.subExecParamGpu + transferOffset; @@ -4487,8 +4918,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } - // 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, @@ -4496,9 +4928,30 @@ const auto& AmdSmiFabricInfoV1(const T& info) exeDevice.exeIndex, exeInfo.wallClockRate}; } else { exeInfo.wallClockRate = 100000; + /* return {ERR_WARN, "GPU %d wallclock rate query returned 0 unexpectedly. Setting to %d instead. Use TB_WALLCLOCK_RATE to customize", exeDevice.exeIndex, exeInfo.wallClockRate}; + */ + } + } + + // Create the persistent, NUMA-pinned worker pool that services this executor across all + // iterations. Sized to the executor's peak intra-executor concurrency. NIC executors + // post work single-threaded and need no pool. + if (exeDevice.exeRank == localRank) { + if (exeDevice.exeType == EXE_CPU) { + int const physNode = System::Get().GetCpuPhysicalNode(exeDevice.exeIndex); + exeInfo.pool.reset(new ThreadPool(std::max(1, exeInfo.totalSubExecs), + [physNode] { numa_run_on_node(physNode); })); + } else if (IsGpuExeType(exeDevice.exeType)) { + int const numThreads = std::max({1, (int)exeInfo.resources.size(), (int)exeInfo.streams.size()}); + int const exeIndex = exeDevice.exeIndex; + int const physNode = (exeNuma >= 0) ? System::Get().GetCpuPhysicalNode(exeNuma) : -1; + exeInfo.pool.reset(new ThreadPool(numThreads, [physNode, exeIndex] { + if (physNode >= 0) numa_run_on_node(physNode); + (void)hipSetDevice(exeIndex); + })); } } @@ -4517,6 +4970,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) int const localRank = GetRank(); bool const verbose = System::Get().IsVerbose(); + // Tear down the persistent worker pool (joins all workers) now that all iterations are done. + exeInfo.pool.reset(); + // Loop over each transfer this executor is involved in for (auto& rss : exeInfo.resources) { Transfer const& t = transfers[rss.transferIdx]; @@ -4573,6 +5029,20 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } + // Deallocate on-device validation buffers (VALIDATE_ON_DEVICE) + for (int iDst = 0; iDst < (int)rss.dstExpectedMem.size(); ++iDst) { + if (rss.dstExpectedMem[iDst]) + ERR_CHECK(DeallocateMemory(t.dsts[iDst].memType, rss.dstExpectedMem[iDst], rss.dstExpectedBytes[iDst])); + if (rss.dstValidateScratch[iDst]) + ERR_CHECK(DeallocateMemory(t.dsts[iDst].memType, rss.dstValidateScratch[iDst], 2 * sizeof(unsigned long long))); + } + for (int iSrc = 0; iSrc < (int)rss.srcExpectedMem.size(); ++iSrc) { + if (rss.srcExpectedMem[iSrc]) + ERR_CHECK(DeallocateMemory(t.srcs[iSrc].memType, rss.srcExpectedMem[iSrc], rss.srcExpectedBytes[iSrc])); + if (rss.srcValidateScratch[iSrc]) + ERR_CHECK(DeallocateMemory(t.srcs[iSrc].memType, rss.srcValidateScratch[iSrc], 2 * sizeof(unsigned long long))); + } + // Destroy HSA signal for DMA executor #if !defined(__NVCC__) if (exeDevice.exeType == EXE_GPU_DMA && (t.exeSubIndex != -1 || cfg.dma.useHsaCopy) && exeDevice.exeRank == localRank) { @@ -4587,19 +5057,17 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Teardown additional requirements for GPU-based executors - if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_DMA || exeDevice.exeType == EXE_GPU_BDMA) - && exeDevice.exeRank == localRank) { + if (IsGpuExeType(exeDevice.exeType) && exeDevice.exeRank == localRank) { for (auto stream : exeInfo.streams) ERR_CHECK(hipStreamDestroy(stream)); - if (cfg.gfx.useHipEvents || cfg.dma.useHipEvents) { - for (auto event : exeInfo.startEvents) - ERR_CHECK(hipEventDestroy(event)); - for (auto event : exeInfo.stopEvents) - ERR_CHECK(hipEventDestroy(event)); - } + for (auto event : exeInfo.startEvents) + ERR_CHECK(hipEventDestroy(event)); + for (auto event : exeInfo.stopEvents) + ERR_CHECK(hipEventDestroy(event)); } - if (exeDevice.exeType == EXE_GPU_GFX && exeDevice.exeRank == localRank) { + if ((exeDevice.exeType == EXE_GPU_GFX || exeDevice.exeType == EXE_GPU_TDM) && + exeDevice.exeRank == localRank) { #if !defined(__NVCC__) MemType memType = MEM_GPU; #else @@ -4655,57 +5123,62 @@ const auto& AmdSmiFabricInfoV1(const T& info) } while (++subIteration != numSubIterations); } - // Execution of a single CPU Transfers - static void ExecuteCpuTransfer(int const iteration, - ConfigOptions const& cfg, - int const exeIndex, - TransferResources& rss) - { - auto cpuStart = std::chrono::high_resolution_clock::now(); - vector childThreads; - - for (auto const& subExecParam : rss.subExecParamCpu) - childThreads.emplace_back(std::thread(CpuReduceKernel, std::cref(subExecParam), cfg.general.numSubIterations)); - - for (auto& subExecThread : childThreads) - subExecThread.join(); - childThreads.clear(); - - 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) { - rss.totalDurationMsec += deltaMsec; - if (cfg.general.recordPerIteration) - rss.perIterMsec.push_back(deltaMsec); - } - } - // Execution of a single CPU executor static ErrResult RunCpuExecutor(int const iteration, ConfigOptions const& cfg, int const exeIndex, ExeInfo& exeInfo) { - numa_run_on_node(exeIndex); - auto cpuStart = std::chrono::high_resolution_clock::now(); + using Clock = std::chrono::high_resolution_clock; + + // Flatten every subexecutor across all of this executor's Transfers into a single task + // list serviced by the persistent, NUMA-pinned worker pool. This collapses the former + // per-transfer + per-subexecutor thread creation into zero per-iteration thread spawns. + struct CpuTask { SubExecParam const* param; int rssIdx; }; + std::vector tasks; + tasks.reserve(exeInfo.totalSubExecs); + for (int r = 0; r < (int)exeInfo.resources.size(); ++r) + for (auto const& p : exeInfo.resources[r].subExecParamCpu) + tasks.push_back({&p, r}); + + int const numTasks = (int)tasks.size(); + std::vector starts(numTasks), stops(numTasks); + + auto cpuStart = Clock::now(); + exeInfo.pool->ParallelFor(numTasks, [&](int i) { + starts[i] = Clock::now(); + CpuReduceKernel(*tasks[i].param, cfg.general.numSubIterations); + stops[i] = Clock::now(); + }); + auto cpuDelta = Clock::now() - cpuStart; - vector asyncTransfers; - for (auto& rss : exeInfo.resources) { - asyncTransfers.emplace_back(std::thread(ExecuteCpuTransfer, - iteration, - std::cref(cfg), - exeIndex, - std::ref(rss))); + if (iteration >= 0) { + // Executor duration: wall-clock span of the whole batch + exeInfo.totalDurationMsec += std::chrono::duration_cast>(cpuDelta).count() + * 1000.0 / cfg.general.numSubIterations; + + // Per-transfer duration: span from the earliest subexecutor start to the latest stop + // among that Transfer's subexecutors (they run concurrently in the pool). + for (int r = 0; r < (int)exeInfo.resources.size(); ++r) { + TransferResources& rss = exeInfo.resources[r]; + auto minStart = Clock::time_point::max(); + auto maxStop = Clock::time_point::min(); + bool any = false; + for (int i = 0; i < numTasks; ++i) { + if (tasks[i].rssIdx != r) continue; + any = true; + minStart = std::min(minStart, starts[i]); + maxStop = std::max(maxStop, stops[i]); + } + double deltaMsec = any + ? std::chrono::duration_cast>(maxStop - minStart).count() + * 1000.0 / cfg.general.numSubIterations + : 0.0; + rss.totalDurationMsec += deltaMsec; + if (cfg.general.recordPerIteration) + rss.perIterMsec.push_back(deltaMsec); + } } - for (auto& asyncTransfer : asyncTransfers) - asyncTransfer.join(); - - 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; } @@ -4813,20 +5286,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) // GFX Executor-related functions //======================================================================================== - // Converts register value to a CU/SM index - static uint32_t GetId(uint32_t hwId) - { -#if defined(__NVCC_) - return hwId; -#else - // Based on instinct-mi200-cdna2-instruction-set-architecture.pdf - int const shId = (hwId >> 12) & 1; - int const cuId = (hwId >> 8) & 15; - int const seId = (hwId >> 13) & 3; - return (shId << 5) + (cuId << 2) + seId; -#endif - } - // Device level timestamp function __device__ int64_t GetTimestamp() { @@ -4959,11 +5418,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (seType == 1 && p.N == 0) return; // Filter by XCC -#if !defined(__NVCC__) - int32_t xccId; - GetXccId(xccId); - if (p.preferredXccId != -1 && xccId != p.preferredXccId) return; -#endif + { uint32_t xccId, cuId; GetXccHwId(xccId, cuId); + if (p.preferredXccId != -1 && (int32_t)xccId != p.preferredXccId) return; } // Collect data information bool hasSrc = p.numSrcs > 0; @@ -5089,8 +5545,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) __threadfence_system(); p.stopCycle = GetTimestamp(); p.startCycle = startCycle; - GetHwId(p.hwId); - GetXccId(p.xccId); + GetXccHwId(p.xccId, p.hwId); } } @@ -5122,11 +5577,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (seType == 1 && p.N == 0) return; // Filter by XCC -#if !defined(__NVCC__) - int32_t xccId; - GetXccId(xccId); - if (p.preferredXccId != -1 && xccId != p.preferredXccId) return; -#endif + { uint32_t xccId, cuId; GetXccHwId(xccId, cuId); + if (p.preferredXccId != -1 && (int32_t)xccId != p.preferredXccId) return; } // Collect data information int32_t const numSrcs = p.numSrcs; @@ -5270,8 +5722,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) __threadfence_system(); p.stopCycle = GetTimestamp(); p.startCycle = startCycle; - GetHwId(p.hwId); - GetXccId(p.xccId); + GetXccHwId(p.xccId, p.hwId); } } @@ -5379,32 +5830,32 @@ const auto& AmdSmiFabricInfoV1(const T& info) #endif // Compute kernel launch parameters - int const numSubExecs = cfg.gfx.useMultiStream ? rss.subExecParamCpu.size() : exeTotalSubExecs; + int const numSubExecs = cfg.general.useMultiStream ? rss.subExecParamCpu.size() : exeTotalSubExecs; int const gridY = CalculateGridY(cfg.gfx.seType, cfg.gfx.blockSize, numSubExecs); dim3 const gridSize(xccDim, gridY, 1); dim3 const blockSize(cfg.gfx.blockSize); auto cpuStart = std::chrono::high_resolution_clock::now(); - SubExecParam* params = cfg.gfx.useMultiStream ? rss.subExecParamGpuPtr : exeSubExecParam; + SubExecParam* params = cfg.general.useMultiStream ? rss.subExecParamGpuPtr : exeSubExecParam; #if defined(__NVCC__) - if (cfg.gfx.useHipEvents) + if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); gpuKernel<<>>(params, cfg.gfx.seType, cfg.gfx.waveOrder, cfg.general.numSubIterations); - if (cfg.gfx.useHipEvents) + if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(stopEvent, stream)); #else hipExtLaunchKernelGGL(gpuKernel, gridSize, blockSize, 0, stream, - cfg.gfx.useHipEvents ? startEvent : NULL, - cfg.gfx.useHipEvents ? stopEvent : NULL, 0, + cfg.general.useHipEvents ? startEvent : NULL, + cfg.general.useHipEvents ? stopEvent : NULL, 0, params, cfg.gfx.seType, cfg.gfx.waveOrder, cfg.general.numSubIterations); #endif ERR_CHECK(hipStreamSynchronize(stream)); // Record this timing if this Transfer is being run in multistream mode - if (cfg.gfx.useMultiStream) { + if (cfg.general.useMultiStream) { if (iteration >= 0) { auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; double cpuDeltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 / cfg.general.numSubIterations; @@ -5430,7 +5881,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } for (int i = 0; i < numSubExecs; i++) { CUs.insert(std::make_pair(subExecParam[i].xccId, - GetId(subExecParam[i].hwId))); + subExecParam[i].hwId)); } rss.perIterCUs.push_back(CUs); } @@ -5450,31 +5901,29 @@ const auto& AmdSmiFabricInfoV1(const T& info) int xccDim = exeInfo.useSubIndices ? exeInfo.numSubIndices : 1; - if (cfg.gfx.useMultiStream) { - // Launch one thread per Transfer in separate streams - vector> asyncTransfers; - for (int i = 0; i < exeInfo.streams.size(); i++) { - asyncTransfers.emplace_back(std::async(std::launch::async, - ExecuteGpuTransfer, - iteration, - exeInfo.totalSubExecs, - exeInfo.subExecParamGpu, - exeInfo.streams[i], - cfg.gfx.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.gfx.useHipEvents ? exeInfo.stopEvents[i] : NULL, - xccDim, - std::cref(cfg), - exeInfo.gfxKernelToUse, - exeInfo.subExecParamHostAccessible, - std::ref(exeInfo.resources[i]))); - } - for (auto& asyncTransfer : asyncTransfers) - ERR_CHECK(asyncTransfer.get()); + if (cfg.general.useMultiStream) { + // Launch one task per Transfer in separate streams on the persistent worker pool + int const numStreams = (int)exeInfo.streams.size(); + std::vector tfrErr(numStreams); + exeInfo.pool->ParallelFor(numStreams, [&](int i) { + tfrErr[i] = ExecuteGpuTransfer(iteration, + exeInfo.totalSubExecs, + exeInfo.subExecParamGpu, + exeInfo.streams[i], + cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, + xccDim, + cfg, + exeInfo.gfxKernelToUse, + exeInfo.subExecParamHostAccessible, + exeInfo.resources[i]); + }); + for (auto& e : tfrErr) ERR_CHECK(e); } else { // Launch all Transfers in one kernel launch (avoid extra thread creation) ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], - cfg.gfx.useHipEvents ? exeInfo.startEvents[0] : NULL, - cfg.gfx.useHipEvents ? exeInfo.stopEvents[0] : NULL, + cfg.general.useHipEvents ? exeInfo.startEvents[0] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[0] : NULL, xccDim, cfg, exeInfo.gfxKernelToUse, exeInfo.subExecParamHostAccessible, exeInfo.resources[0]); } @@ -5485,7 +5934,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Determine executor timing // - Use HIP event timing if enabled and not using multi-stream // - Otherwise, Use CPU timing - if (cfg.gfx.useHipEvents && !cfg.gfx.useMultiStream) { + if (cfg.general.useHipEvents && !cfg.general.useMultiStream) { float gpuDeltaMsec; ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); gpuDeltaMsec /= cfg.general.numSubIterations; @@ -5498,7 +5947,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) // If Transfers were combined into a single launch, figure out per-Transfer timing // Determine timing for each of the individual transfers that were part of this launch - if (!cfg.gfx.useMultiStream) { + if (!cfg.general.useMultiStream) { std::vector subExecParamHost; SubExecParam const* subExecParam = exeInfo.subExecParamGpu; if (!exeInfo.subExecParamHostAccessible) { @@ -5520,7 +5969,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) maxStopCycle = std::max(maxStopCycle, subExecParam[subExecIdx].stopCycle); if (cfg.general.recordPerIteration) { CUs.insert(std::make_pair(subExecParam[subExecIdx].xccId, - GetId(subExecParam[subExecIdx].hwId))); + subExecParam[subExecIdx].hwId)); } } } @@ -5559,7 +6008,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) size_t const initOffset = cfg.data.byteOffset / sizeof(float); float* const src = resources.srcMem[0] + initOffset; if (!useSubIndices && !cfg.dma.useHsaCopy) { - if (cfg.dma.useHipEvents) + if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); // Force the use of SDMA engine if possible @@ -5583,7 +6032,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } while (++subIterations != cfg.general.numSubIterations); - if (cfg.dma.useHipEvents) + if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(stopEvent, stream)); ERR_CHECK(hipStreamSynchronize(stream)); } else { @@ -5620,7 +6069,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (iteration >= 0) { double deltaMsec = cpuDeltaMsec; - if (!useSubIndices && !cfg.dma.useHsaCopy && cfg.dma.useHipEvents) { + if (!useSubIndices && !cfg.dma.useHsaCopy && cfg.general.useHipEvents) { float gpuDeltaMsec; ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations; @@ -5641,22 +6090,19 @@ const auto& AmdSmiFabricInfoV1(const T& info) auto cpuStart = std::chrono::high_resolution_clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); - vector> asyncTransfers; - for (int i = 0; i < exeInfo.resources.size(); i++) { - asyncTransfers.emplace_back(std::async(std::launch::async, - ExecuteDmaTransfer, - iteration, - exeInfo.useSubIndices, - exeIndex, - exeInfo.streams[i], - cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL, - std::cref(cfg), - std::ref(exeInfo.resources[i]))); - } - - for (auto& asyncTransfer : asyncTransfers) - ERR_CHECK(asyncTransfer.get()); + int const numTransfers = (int)exeInfo.resources.size(); + std::vector tfrErr(numTransfers); + exeInfo.pool->ParallelFor(numTransfers, [&](int i) { + tfrErr[i] = ExecuteDmaTransfer(iteration, + exeInfo.useSubIndices, + exeIndex, + exeInfo.streams[i], + cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, + cfg, + exeInfo.resources[i]); + }); + for (auto& e : tfrErr) ERR_CHECK(e); auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 / cfg.general.numSubIterations; @@ -5682,7 +6128,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) ERR_CHECK(hipSetDevice(exeIndex)); int subIterations = 0; - if (cfg.dma.useHipEvents) + if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(startEvent, stream)); [[maybe_unused]] size_t failIdx = 0; @@ -5699,7 +6145,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) stream)); } while (++subIterations != cfg.general.numSubIterations); - if (cfg.dma.useHipEvents) + if (cfg.general.useHipEvents) ERR_CHECK(hipEventRecord(stopEvent, stream)); ERR_CHECK(hipStreamSynchronize(stream)); @@ -5708,7 +6154,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (iteration >= 0) { double deltaMsec = cpuDeltaMsec; - if (cfg.dma.useHipEvents) { + if (cfg.general.useHipEvents) { float gpuDeltaMsec; ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, startEvent, stopEvent)); deltaMsec = gpuDeltaMsec / cfg.general.numSubIterations; @@ -5728,21 +6174,18 @@ const auto& AmdSmiFabricInfoV1(const T& info) auto cpuStart = std::chrono::high_resolution_clock::now(); ERR_CHECK(hipSetDevice(exeIndex)); - vector> asyncTransfers; - for (int i = 0; i < exeInfo.resources.size(); i++) { - asyncTransfers.emplace_back(std::async(std::launch::async, - ExecuteBatchDmaTransfer, - iteration, - exeIndex, - exeInfo.streams[i], - cfg.dma.useHipEvents ? exeInfo.startEvents[i] : NULL, - cfg.dma.useHipEvents ? exeInfo.stopEvents[i] : NULL, - std::cref(cfg), - std::ref(exeInfo.resources[i]))); - } - - for (auto& asyncTransfer : asyncTransfers) - ERR_CHECK(asyncTransfer.get()); + int const numTransfers = (int)exeInfo.resources.size(); + std::vector tfrErr(numTransfers); + exeInfo.pool->ParallelFor(numTransfers, [&](int i) { + tfrErr[i] = ExecuteBatchDmaTransfer(iteration, + exeIndex, + exeInfo.streams[i], + cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, + cfg, + exeInfo.resources[i]); + }); + for (auto& e : tfrErr) ERR_CHECK(e); auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; double deltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 / cfg.general.numSubIterations; @@ -5752,6 +6195,209 @@ const auto& AmdSmiFabricInfoV1(const T& info) } #endif // BMA_EXEC_ENABLED +// TDM Executor-related functions +//======================================================================================== +#if TDM_SUPPORTED + __global__ void GpuTdmKernel(SubExecParam* params, + uint32_t ldsBytes, + int numSubIterations) + { + 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). + SubExecParam& p = params[blockIdx.x]; + if (p.N == 0) return; + + float const* __restrict__ src = (float const*)p.src[0]; + float* __restrict__ dst = (float*)p.dst[0]; + size_t const sizeBytes = p.N * sizeof(float); + + int subIterations = 0; + while (1) { + tdm::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes); + __syncthreads(); // Wait for all warps to finish + if (++subIterations == numSubIterations) break; + } + + if (shouldRecordTiming) { + __threadfence_system(); + p.stopCycle = GetTimestamp(); + p.startCycle = startCycle; + GetXccHwId(p.xccId, p.hwId); + } + } +#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*, uint32_t, int) {} +#endif // TDM_SUPPORTED + + static ErrResult ExecuteTdmTransfer(int const iteration, + int const exeTotalSubExecs, + SubExecParam* exeSubExecParam, + hipStream_t const stream, + hipEvent_t const startEvent, + hipEvent_t const stopEvent, + ConfigOptions const& cfg, + bool const subExecParamHostAccessible, + uint32_t const ldsBytes, + TransferResources& rss) + { + // Compute kernel launch parameters + int const numSubExecs = cfg.general.useMultiStream ? rss.subExecParamCpu.size() : exeTotalSubExecs; + dim3 const gridSize(numSubExecs); + dim3 const blockSize(cfg.tdm.blockSize); + SubExecParam* params = cfg.general.useMultiStream ? rss.subExecParamGpuPtr : exeSubExecParam; + + auto cpuStart = std::chrono::high_resolution_clock::now(); + +#if defined(__NVCC__) + if (cfg.general.useHipEvents) + ERR_CHECK(hipEventRecord(startEvent, stream)); + GpuTdmKernel<<>>(params, + ldsBytes, + cfg.general.numSubIterations); + if (cfg.general.useHipEvents) + ERR_CHECK(hipEventRecord(stopEvent, stream)); +#else + hipExtLaunchKernelGGL(GpuTdmKernel, gridSize, blockSize, (int)ldsBytes, stream, + startEvent, stopEvent, 0, + params, ldsBytes, cfg.general.numSubIterations); +#endif + ERR_CHECK(hipStreamSynchronize(stream)); + + // Record this timing if this Transfer is being run in multistream mode + if (cfg.general.useMultiStream) { + if (iteration >= 0) { + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + double cpuDeltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 / cfg.general.numSubIterations; + + double deltaMsec = cpuDeltaMsec; + if (startEvent != NULL) { + 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; + std::vector subExecParamHost; + SubExecParam const* subExecParam = rss.subExecParamGpuPtr; + if (!subExecParamHostAccessible) { + subExecParamHost.resize(numSubExecs); + ERR_CHECK(hipMemcpy(subExecParamHost.data(), rss.subExecParamGpuPtr, + numSubExecs * sizeof(SubExecParam), hipMemcpyDefault)); + subExecParam = subExecParamHost.data(); + } + for (int i = 0; i < numSubExecs; i++) { + CUs.insert(std::make_pair(subExecParam[i].xccId, + subExecParam[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)); + + if (cfg.general.useMultiStream && exeInfo.streams.size() > 1 ) { + // Launch one task per Transfer in separate streams on the persistent worker pool + int const numStreams = (int)exeInfo.streams.size(); + std::vector tfrErr(numStreams); + exeInfo.pool->ParallelFor(numStreams, [&](int i) { + tfrErr[i] = ExecuteTdmTransfer(iteration, + exeInfo.totalSubExecs, + exeInfo.subExecParamGpu, + exeInfo.streams[i], + cfg.general.useHipEvents ? exeInfo.startEvents[i] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[i] : NULL, + cfg, + exeInfo.subExecParamHostAccessible, + exeInfo.ldsBytesActual, + exeInfo.resources[i]); + }); + for (auto& e : tfrErr) ERR_CHECK(e); + } else { + // Launch all Transfers in one kernel launch (avoid extra thread creation) + ExecuteTdmTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0], + cfg.general.useHipEvents ? exeInfo.startEvents[0] : NULL, + cfg.general.useHipEvents ? exeInfo.stopEvents[0] : NULL, + cfg, exeInfo.subExecParamHostAccessible, exeInfo.ldsBytesActual, exeInfo.resources[0]); + } + + if (iteration >= 0) { + // Determine executor timing + // - Use HIP event timing if enabled and not using multi-stream + // - Otherwise, Use CPU timing + if (cfg.general.useHipEvents && !cfg.general.useMultiStream) { + float gpuDeltaMsec; + ERR_CHECK(hipEventElapsedTime(&gpuDeltaMsec, exeInfo.startEvents[0], exeInfo.stopEvents[0])); + gpuDeltaMsec /= cfg.general.numSubIterations; + exeInfo.totalDurationMsec += gpuDeltaMsec; + } else { + auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart; + double cpuDeltaMsec = std::chrono::duration_cast>(cpuDelta).count() * 1000.0 + / cfg.general.numSubIterations; + exeInfo.totalDurationMsec += cpuDeltaMsec; + } + + // If Transfers were combined into a single launch, figure out per-Transfer timing + // Determine timing for each of the individual transfers that were part of this launch + if (!cfg.general.useMultiStream) { + std::vector subExecParamHost; + SubExecParam const* subExecParam = exeInfo.subExecParamGpu; + if (!exeInfo.subExecParamHostAccessible) { + subExecParamHost.resize(exeInfo.totalSubExecs); + ERR_CHECK(hipMemcpy(subExecParamHost.data(), exeInfo.subExecParamGpu, + exeInfo.totalSubExecs * sizeof(SubExecParam), hipMemcpyDefault)); + subExecParam = subExecParamHost.data(); + } + + for (int i = 0; i < exeInfo.resources.size(); i++) { + TransferResources& rss = exeInfo.resources[i]; + int64_t minStartCycle = std::numeric_limits::max(); + int64_t maxStopCycle = std::numeric_limits::min(); + std::set> CUs; + + for (auto subExecIdx : rss.subExecIdx) { + if (exeInfo.subExecParamCpu[subExecIdx].N != 0) { + minStartCycle = std::min(minStartCycle, subExecParam[subExecIdx].startCycle); + maxStopCycle = std::max(maxStopCycle, subExecParam[subExecIdx].stopCycle); + if (cfg.general.recordPerIteration) { + CUs.insert(std::make_pair(subExecParam[subExecIdx].xccId, + subExecParam[subExecIdx].hwId)); + } + } + } + + double deltaMsec = (maxStopCycle - minStartCycle) / (double)(exeInfo.wallClockRate); + deltaMsec /= cfg.general.numSubIterations; + rss.totalDurationMsec += deltaMsec; + if (cfg.general.recordPerIteration) { + rss.perIterMsec.push_back(deltaMsec); + rss.perIterCUs.push_back(CUs); + } + } + } + } + return ERR_NONE; + } + // Executor-related functions //======================================================================================== static ErrResult RunExecutor(int const iteration, @@ -5761,8 +6407,9 @@ const auto& AmdSmiFabricInfoV1(const T& info) { 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_DMA: return RunDmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_GFX: return RunGpuExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); + case EXE_GPU_TDM: return RunTdmExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); case EXE_NIC: return RunNicExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); #ifdef BMA_EXEC_ENABLED case EXE_GPU_BDMA: return RunBmaExecutor(iteration, cfg, exeDevice.exeIndex, exeInfo); @@ -5975,6 +6622,63 @@ const auto& AmdSmiFabricInfoV1(const T& info) ERR_APPEND(hipMemcpy(resource->srcMem[srcIdx] + initOffset, srcReference[srcIdx].data(), resource->numBytes, hipMemcpyDefault), errResults); ERR_APPEND(hipDeviceSynchronize(), errResults); + + // Optionally validate source memory right after preparation + if (cfg.data.validateSource) { + if (cfg.data.validateOnDevice && IsGpuMemType(t.srcs[srcIdx].memType)) { + resource->srcExpectedMem.resize(resource->srcMem.size(), nullptr); + resource->srcExpectedBytes.resize(resource->srcMem.size(), 0); + resource->srcValidateScratch.resize(resource->srcMem.size(), nullptr); + ERR_APPEND(AllocateMemory(t.srcs[srcIdx], resource->numBytes, + (void**)&resource->srcExpectedMem[srcIdx], + &resource->srcExpectedBytes[srcIdx]), errResults); + ERR_APPEND(AllocateMemory(t.srcs[srcIdx], 2 * sizeof(unsigned long long), + (void**)&resource->srcValidateScratch[srcIdx]), errResults); + ERR_APPEND(hipMemcpy(resource->srcExpectedMem[srcIdx], srcReference[srcIdx].data(), + resource->numBytes, hipMemcpyDefault), errResults); + ERR_APPEND(hipDeviceSynchronize(), errResults); + unsigned long long count = 0, firstOff = 0; + ERR_APPEND(ValidateBufferOnDevice(t.srcs[srcIdx].memIndex, + resource->srcMem[srcIdx] + initOffset, + resource->srcExpectedMem[srcIdx], resource->numBytes, + resource->srcValidateScratch[srcIdx], &count, &firstOff), errResults); + if (count) + ERR_APPEND((ErrResult{ERR_FATAL, "Transfer %d: Source %d mismatch after prep: %llu mismatched bytes, first at offset %llu", + resource->transferIdx, srcIdx, count, firstOff}), errResults); + } else { + size_t const N = resource->numBytes / sizeof(float); + vector tmp(N); + ERR_APPEND(hipMemcpy(tmp.data(), resource->srcMem[srcIdx] + initOffset, resource->numBytes, hipMemcpyDefault), errResults); + ERR_APPEND(hipDeviceSynchronize(), errResults); + if (memcmp(tmp.data(), srcReference[srcIdx].data(), resource->numBytes)) + ERR_APPEND((ErrResult{ERR_FATAL, "Transfer %d: Source %d mismatch after prep", + resource->transferIdx, srcIdx}), errResults); + } + } + } + } + } + + // Pre-upload expected destination values to device for on-device validation + if (validateEnabled && cfg.data.validateOnDevice) { + for (auto resource : transferResources) { + Transfer const& t = transfers[resource->transferIdx]; + float const* dstExpected = dstReference[t.srcs.size()].data(); + resource->dstExpectedMem.resize(resource->dstMem.size(), nullptr); + resource->dstExpectedBytes.resize(resource->dstMem.size(), 0); + resource->dstValidateScratch.resize(resource->dstMem.size(), nullptr); + for (int dstIdx = 0; dstIdx < (int)resource->dstMem.size(); dstIdx++) { + if (t.dsts[dstIdx].memRank != localRank) continue; + if (!IsGpuMemType(t.dsts[dstIdx].memType)) continue; + ERR_APPEND(hipSetDevice(t.dsts[dstIdx].memIndex), errResults); + ERR_APPEND(AllocateMemory(t.dsts[dstIdx], resource->numBytes, + (void**)&resource->dstExpectedMem[dstIdx], + &resource->dstExpectedBytes[dstIdx]), errResults); + ERR_APPEND(AllocateMemory(t.dsts[dstIdx], 2 * sizeof(unsigned long long), + (void**)&resource->dstValidateScratch[dstIdx]), errResults); + ERR_APPEND(hipMemcpy(resource->dstExpectedMem[dstIdx], dstExpected, + resource->numBytes, hipMemcpyDefault), errResults); + ERR_APPEND(hipDeviceSynchronize(), errResults); } } } @@ -6031,6 +6735,18 @@ const auto& AmdSmiFabricInfoV1(const T& info) } // Perform iterations + // Persistent pool that dispatches the local executors concurrently each iteration, + // replacing per-iteration std::async. Each RunExecutor still sets its own device/NUMA. + // ExeInfo pointers are resolved up-front so the parallel lambda never touches the map + // (std::map::operator[] is non-const and not safe to call concurrently). + std::vector localExeInfos; + localExeInfos.reserve(localExecutors.size()); + for (auto const& exeDevice : localExecutors) + localExeInfos.push_back(&executorMap[exeDevice]); + + ThreadPool executorPool((int)localExecutors.size()); + std::vector exeErrors(localExecutors.size()); + size_t numTimedIterations = 0; double totalCpuTimeSec = 0.0; for (int iteration = -cfg.general.numWarmups; ; iteration++) { @@ -6048,19 +6764,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Start CPU timing for this iteration auto cpuStart = std::chrono::high_resolution_clock::now(); - // Execute all Transfers in parallel - std::vector> asyncExecutors; - for (auto const& exeDevice : localExecutors) { - asyncExecutors.emplace_back(std::async(std::launch::async, RunExecutor, - iteration, - std::cref(cfg), - std::cref(exeDevice), - std::ref(executorMap[exeDevice]))); - } + // Execute all local executors in parallel on the persistent executor pool + executorPool.ParallelFor((int)localExecutors.size(), [&](int i) { + exeErrors[i] = RunExecutor(iteration, cfg, localExecutors[i], *localExeInfos[i]); + }); - // Wait for all threads to finish - for (auto& asyncExecutor : asyncExecutors) { - ERR_APPEND(asyncExecutor.get(), errResults); + // Collect any errors reported by the executors + for (auto& exeErr : exeErrors) { + ERR_APPEND(exeErr, errResults); } // Wait for all ranks to finish @@ -6073,6 +6784,23 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (cfg.data.alwaysValidate > 0) { ERR_APPEND(ValidateAllTransfers(cfg, transfers, transferResources, dstReference, outputBuffer), errResults); + + // Clear destination memory after validation so each iteration starts from a known-zero state + size_t const initOffset = cfg.data.byteOffset / sizeof(float); + for (auto rss : transferResources) { + Transfer const& t = transfers[rss->transferIdx]; + for (int dstIdx = 0; dstIdx < (int)rss->dstMem.size(); dstIdx++) { + if (t.dsts[dstIdx].memRank != localRank) continue; + float* dstPtr = rss->dstMem[dstIdx] + initOffset; + if (IsCpuMemType(t.dsts[dstIdx].memType)) { + memset(dstPtr, 0, rss->numBytes); + } else if (IsGpuMemType(t.dsts[dstIdx].memType)) { + ERR_APPEND(hipSetDevice(t.dsts[dstIdx].memIndex), errResults); + ERR_APPEND(hipMemset(dstPtr, 0, rss->numBytes), errResults); + ERR_APPEND(hipDeviceSynchronize(), errResults); + } + } + } } if (iteration >= 0) { @@ -6112,7 +6840,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) for (auto rss : transferResources) { int transferIdx = rss->transferIdx; Transfer const& t = transfers[transferIdx]; - size_t N = t.numBytes / sizeof(float); float const* expected = dstReference[t.srcs.size()].data(); bool transferOk = true; bool anyLocalDst = false; @@ -6124,6 +6851,19 @@ const auto& AmdSmiFabricInfoV1(const T& info) continue; } anyLocalDst = true; + if (cfg.data.validateOnDevice && IsGpuMemType(t.dsts[dstIdx].memType)) { + unsigned long long count = 0, firstOff = 0; + (void)ValidateBufferOnDevice(t.dsts[dstIdx].memIndex, rss->dstMem[dstIdx] + initOffset, + rss->dstExpectedMem[dstIdx], t.numBytes, + rss->dstValidateScratch[dstIdx], &count, &firstOff); + if (count == 0) { + System::Get().Log(" DST[%d]=PASS", dstIdx); + } else { + System::Get().Log(" DST[%d]=FAIL(%llu bytes, first @ %llu)", dstIdx, count, firstOff); + transferOk = false; + } + continue; + } float* output; if (IsCpuMemType(t.dsts[dstIdx].memType) || cfg.data.validateDirect) { output = rss->dstMem[dstIdx] + initOffset; @@ -6137,15 +6877,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) if (memcmp(output, expected, t.numBytes) == 0) { System::Get().Log(" DST[%d]=PASS", dstIdx); } else { - size_t firstErr = 0; - for (; firstErr < N; firstErr++) - if (output[firstErr] != expected[firstErr]) break; - if (firstErr < N) - System::Get().Log(" DST[%d]=FAIL(first mismatch idx=%zu exp=%.5f got=%.5f)", - dstIdx, firstErr, expected[firstErr], output[firstErr]); - else - System::Get().Log(" DST[%d]=FAIL(bitwise mismatch, no float-level diff found)", - dstIdx); + std::string ranges = ByteMismatchSummary(output, expected, t.numBytes); + System::Get().Log(" DST[%d]=FAIL(bytes %s)", dstIdx, ranges.c_str()); transferOk = false; } } @@ -6480,7 +7213,8 @@ const auto& AmdSmiFabricInfoV1(const T& info) return result; } else if (wc.exe.exeSubIndices[0] == -2) { switch (wc.exe.exeType) { - case EXE_CPU: + case EXE_CPU: case EXE_GPU_TDM: + // These Executors do not support subindices wc.exe.exeSubIndices[0] = -1; result |= RecursiveWildcardTransferExpansion(wc, baseRankIndex, numBytes, numSubExecs, transfers); wc.exe.exeSubIndices[0] = -2; @@ -6658,9 +7392,13 @@ const auto& AmdSmiFabricInfoV1(const T& info) // TB_SINGLE_LOG = Only rank 0 will produce output (useful if spawning multi-node socket) // TB_DUMP_CFG_FILE = Config file to dump executed Transfers // TB_PAUSE = Insert a pause for debug attachment + // TB_SHOW_ALL_NUMA = Expose all CPU NUMA nodes (default skips nodes with no cores) + // TB_SEND_USLEEP = microseconds to sleep after each socket SendData (default 0) - verbose = getenv("TB_VERBOSE") ? atoi(getenv("TB_VERBOSE")) : 0; + verbose = getenv("TB_VERBOSE") ? atoi(getenv("TB_VERBOSE")) : 0; + sendUsleep = getenv("TB_SEND_USLEEP") ? atoi(getenv("TB_SEND_USLEEP")) : 0; bool singleLog = getenv("TB_SINGLE_LOG") ? atoi(getenv("TB_SINGLE_LOG")) : 0; + showAllNuma = getenv("TB_SHOW_ALL_NUMA") ? atoi(getenv("TB_SHOW_ALL_NUMA")) : 0; char* dumpCfgFilename = getenv("TB_DUMP_CFG_FILE"); if (dumpCfgFilename) { @@ -6701,8 +7439,15 @@ const auto& AmdSmiFabricInfoV1(const T& info) Log("[INFO] Running in single node mode\n"); } + // Build the CPU NUMA remapping before collecting topology so that all + // subsequent CPU indexing (agents, executors, memory) is consistent. + BuildCpuNumaMap(); + // Collect topology and distribute across all ranks CollectTopology(); + if (verbose) { + Log("[INFO] Finished topology exchange\n"); + } } System::~System() @@ -6874,6 +7619,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } sockets.resize(numRanks, -1); + socketEpoch.resize(numRanks, 0); // Rank 0 acts as server for others to connect to int opt = 1; @@ -6937,9 +7683,20 @@ const auto& AmdSmiFabricInfoV1(const T& info) exit(1); } + // Disable Nagle's algorithm to prevent small-message coalescing delays + setsockopt(clientSocket, IPPROTO_TCP, TCP_NODELAY, &opt, sizeof(opt)); + // Receive rank ID from client - int clientRank; - recv(clientSocket, (char*)&clientRank, sizeof(clientRank), 0); + int clientRank = -1; + size_t totalRecv = 0; + while (totalRecv < sizeof(clientRank)) { + auto r = recv(clientSocket, (char*)&clientRank + totalRecv, sizeof(clientRank) - totalRecv, 0); + if (r <= 0) { + Log("[ERROR] Failed to receive rank ID from connecting client\n"); + exit(1); + } + totalRecv += r; + } if (clientRank < 0 || clientRank >= numRanks) { close(clientSocket); @@ -6951,6 +7708,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) } sockets[clientRank] = clientSocket; } + Log("[INFO] %d other rank(s) have connected\n", numRanks - 1); } else { // All other ranks connect to rank 0 int sock = socket(AF_INET, SOCK_STREAM, IPPROTO_TCP); @@ -6991,8 +7749,19 @@ const auto& AmdSmiFabricInfoV1(const T& info) exit(1); } + // Disable Nagle's algorithm to prevent small-message coalescing delays + setsockopt(sock, IPPROTO_TCP, TCP_NODELAY, &opt, sizeof(opt)); + // Send local rank to the server - send(sock, (char*)&rank, sizeof(rank), 0); + size_t totalSent = 0; + while (totalSent < sizeof(rank)) { + auto s = send(sock, (char*)&rank + totalSent, sizeof(rank) - totalSent, MSG_NOSIGNAL); + if (s <= 0) { + Log("[ERROR] Failed to send rank ID to master\n"); + exit(1); + } + totalSent += s; + } sockets[0] = sock; } @@ -7077,7 +7846,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) fprintf(dumpCfgFile, "\n"); } - void System::Barrier() + void System::Barrier() const { #ifdef MPI_COMM_ENABLED if (commMode == COMM_MPI) { @@ -7122,16 +7891,22 @@ const auto& AmdSmiFabricInfoV1(const T& info) } auto sock = sockets[dstRank]; - // Send data + if (verbose) + Log("[SOCK] SendData #%d: rank %d -> rank %d, %zu bytes (usleep=%dus)\n", + socketEpoch[dstRank], rank, dstRank, numBytes, sendUsleep); + socketEpoch[dstRank]++; + size_t totalSent = 0; while (totalSent < numBytes) { - auto sent = send(sock, (char*)sendData + totalSent, numBytes - totalSent, 0); + auto sent = send(sock, (char*)sendData + totalSent, numBytes - totalSent, MSG_NOSIGNAL); if (sent == -1) { Log("[ERROR] Send failed (rank %d to rank %d)\n", rank, dstRank); exit(1); } totalSent += sent; } + + if (sendUsleep > 0) usleep(sendUsleep); } } @@ -7151,6 +7926,12 @@ const auto& AmdSmiFabricInfoV1(const T& info) } auto sock = sockets[srcRank]; + + if (verbose) + Log("[SOCK] RecvData #%d: rank %d <- rank %d, %zu bytes\n", + socketEpoch[srcRank], rank, srcRank, numBytes); + socketEpoch[srcRank]++; + size_t totalRecv = 0; while (totalRecv < numBytes) { auto recvd = recv(sock, (char*)recvData + totalRecv, numBytes - totalRecv, 0); @@ -7323,6 +8104,83 @@ const auto& AmdSmiFabricInfoV1(const T& info) #endif } + // Build the logical<->physical CPU NUMA node mapping. + // Default: skip NUMA nodes that have no CPU cores. + // TB_SHOW_ALL_NUMA=1: reproduce legacy behavior (logical index == physical node). + void System::BuildCpuNumaMap() + { + cpuNumaMap.clear(); + cpuNumaRevMap.clear(); + + if (showAllNuma) { + // Legacy behavior: expose every configured node with identity mapping + int numConfigured = numa_num_configured_nodes(); + for (int node = 0; node < numConfigured; node++) + cpuNumaMap.push_back(node); + } else { + // Skip NUMA nodes that have no CPU cores + int numConfiguredCpus = numa_num_configured_cpus(); + for (int node = 0; node <= numa_max_node(); node++) { + int coreCount = 0; + for (int cpu = 0; cpu < numConfiguredCpus; cpu++) + if (numa_node_of_cpu(cpu) == node) coreCount++; + if (coreCount > 0) + cpuNumaMap.push_back(node); + } + // Safety fallback: if no node reported cores, fall back to legacy enumeration + if (cpuNumaMap.empty()) { + int numConfigured = numa_num_configured_nodes(); + for (int node = 0; node < numConfigured; node++) + cpuNumaMap.push_back(node); + } + } + + for (int logical = 0; logical < (int)cpuNumaMap.size(); logical++) + cpuNumaRevMap[cpuNumaMap[logical]] = logical; + + if (verbose) { + std::string mapStr; + for (int logical = 0; logical < (int)cpuNumaMap.size(); logical++) + mapStr += " " + std::to_string(logical) + "->" + std::to_string(cpuNumaMap[logical]); + Log("[INFO] Rank %03d: CPU NUMA map (logical->physical)%s%s\n", rank, mapStr.c_str(), + showAllNuma ? " [TB_SHOW_ALL_NUMA]" : ""); + } + } + + int System::GetCpuPhysicalNode(int logicalIdx) const + { + if (logicalIdx < 0 || logicalIdx >= (int)cpuNumaMap.size()) + return logicalIdx; + return cpuNumaMap[logicalIdx]; + } + + int System::GetCpuLogicalNode(int physicalNode) const + { + auto it = cpuNumaRevMap.find(physicalNode); + return (it == cpuNumaRevMap.end()) ? -1 : it->second; + } + + int System::GetClosestLogicalCpu(int physicalNode) const + { + if (physicalNode < 0 || cpuNumaMap.empty()) return -1; + + // Exact match: the physical node is itself exposed + int logical = GetCpuLogicalNode(physicalNode); + if (logical >= 0) return logical; + + // Otherwise pick the exposed node with the smallest NUMA distance + int bestLogical = -1; + int bestDist = std::numeric_limits::max(); + for (int i = 0; i < (int)cpuNumaMap.size(); i++) { + int dist = numa_distance(physicalNode, cpuNumaMap[i]); + if (dist < bestDist) { + bestDist = dist; + bestLogical = i; + } + } + return bestLogical; + } + void System::GetRankTopology(RankTopology& topo) { // Clear topology structure first @@ -7342,19 +8200,36 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Collect Pod membership CollectPodMembership(topo.ppodId, topo.vpodId); - // CPU Executor - int numCpus = numa_num_configured_nodes(); + if (verbose) { + if (topo.vpodId == -1) { + Log("[INFO] Rank %03d: No pod membership detected\n", rank); + } else { + auto* p = (unsigned char*)topo.ppodId; + char ppodUuid[37]; + snprintf(ppodUuid, sizeof(ppodUuid), + "%02x%02x%02x%02x-%02x%02x-%02x%02x-%02x%02x-%02x%02x%02x%02x%02x%02x", + p[0],p[1],p[2],p[3], p[4],p[5], p[6],p[7], + p[8],p[9], p[10],p[11],p[12],p[13],p[14],p[15]); + Log("[INFO] Rank %03d: ppod_id=%s vpod_id=%lld\n", rank, ppodUuid, (long long)topo.vpodId); + } + } + + // CPU Executor (indexed by logical CPU NUMA index; core-less nodes may be skipped) + int numCpus = static_cast(cpuNumaMap.size()); topo.numExecutors[EXE_CPU] = numCpus; std::string cpuName = GetCpuName(); for (int exeIndex = 0; exeIndex < numCpus; exeIndex++) { topo.numExecutorSubIndices[{EXE_CPU, exeIndex}] = 0; + topo.numSubExecutors[{EXE_CPU, exeIndex}] = 0; topo.executorName[{EXE_CPU, exeIndex}] = cpuName; } for (int cpuCore = 0; cpuCore < numa_num_configured_cpus(); cpuCore++) { - topo.numSubExecutors[{EXE_CPU, numa_node_of_cpu(cpuCore)}]++; + int logical = GetCpuLogicalNode(numa_node_of_cpu(cpuCore)); + if (logical >= 0) + topo.numSubExecutors[{EXE_CPU, logical}]++; } if (verbose) { @@ -7369,9 +8244,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) int numGpus = 0; hipError_t status = hipGetDeviceCount(&numGpus); if (status != hipSuccess) numGpus = 0; - topo.numExecutors[EXE_GPU_GFX] = numGpus; - topo.numExecutors[EXE_GPU_DMA] = numGpus; + topo.numExecutors[EXE_GPU_GFX] = numGpus; + topo.numExecutors[EXE_GPU_DMA] = numGpus; topo.numExecutors[EXE_GPU_BDMA] = numGpus; + topo.numExecutors[EXE_GPU_TDM] = numGpus; std::vector gpuArchNames(numGpus); @@ -7392,9 +8268,10 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::string fullName = props.gcnArchName; gpuArchNames[exeIndex] = fullName.substr(0, fullName.find(':')); } - topo.executorName[{EXE_GPU_GFX, exeIndex}] = gpuName; - topo.executorName[{EXE_GPU_DMA, exeIndex}] = gpuName; + topo.executorName[{EXE_GPU_GFX, exeIndex}] = gpuName; + topo.executorName[{EXE_GPU_DMA, exeIndex}] = gpuName; topo.executorName[{EXE_GPU_BDMA, exeIndex}] = gpuName; + topo.executorName[{EXE_GPU_TDM, exeIndex}] = gpuName; #if !defined(__NVCC__) hsa_agent_t gpuAgent = gpuAgents[exeIndex]; @@ -7421,15 +8298,18 @@ const auto& AmdSmiFabricInfoV1(const T& info) } } #endif - topo.numExecutorSubIndices[{EXE_GPU_GFX, exeIndex}] = numXccs; - topo.numExecutorSubIndices[{EXE_GPU_DMA, exeIndex}] = numDmaEngines; + topo.numExecutorSubIndices[{EXE_GPU_GFX, exeIndex}] = numXccs; + topo.numExecutorSubIndices[{EXE_GPU_DMA, exeIndex}] = numDmaEngines; topo.numExecutorSubIndices[{EXE_GPU_BDMA, exeIndex}] = 0; - topo.numSubExecutors[{EXE_GPU_GFX, exeIndex}] = numDeviceCUs; - topo.numSubExecutors[{EXE_GPU_DMA, exeIndex}] = 1; + topo.numExecutorSubIndices[{EXE_GPU_TDM, exeIndex}] = 0; + + topo.numSubExecutors[{EXE_GPU_GFX, exeIndex}] = numDeviceCUs; + topo.numSubExecutors[{EXE_GPU_DMA, exeIndex}] = 1; topo.numSubExecutors[{EXE_GPU_BDMA, exeIndex}] = numDmaEngines; + topo.numSubExecutors[{EXE_GPU_TDM, exeIndex}] = numDeviceCUs; + topo.closestCpuNumaToGpu[exeIndex] = closestNuma; topo.closestNicsToGpu[exeIndex] = {}; - } // NIC Executor @@ -7438,7 +8318,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) { numNics = GetIbvDeviceList().size(); for (int exeIndex = 0; exeIndex < numNics; exeIndex++) { - topo.closestCpuNumaToNic[exeIndex] = GetIbvDeviceList()[exeIndex].numaNode; + // Report the closest CPU NUMA as a logical index (matching CPU executor indices). + // If the NIC's physical node is not exposed (e.g. core-less), fall back to the + // nearest exposed node by NUMA distance so the value stays a valid CPU index. + int nicPhysNode = GetIbvDeviceList()[exeIndex].numaNode; + topo.closestCpuNumaToNic[exeIndex] = GetClosestLogicalCpu(nicPhysNode); topo.executorName[{EXE_NIC, exeIndex}] = GetIbvDeviceList()[exeIndex].name; topo.nicIsActive[exeIndex] = GetIbvDeviceList()[exeIndex].hasActivePort; if (verbose) { @@ -7482,19 +8366,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) 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; @@ -7503,34 +8374,31 @@ const auto& AmdSmiFabricInfoV1(const T& info) // 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 + // The following will only use distance between PCIe domains // to determine the closest NIC to GPU if the PCIe tree approach fails - if (closestIdx < 0) { + if (closestNicIdxs.empty()) { #ifdef VERBS_DEBUG - Log("[WARN] Falling back to PCIe bus ID distance to determine proximity\n"); + Log("[WARN] Falling back to PCIe domain 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) { + // Use ibvAddressList rather than the raw device list: it is already blanked out + // for NICs without an active port, so this stays consistent with the tree path + // above and never maps a GPU to a NIC that cannot execute a Transfer + if (ibvAddressList[nicIndex] != "") { + int distance = GetDomainDistance(hipPciBusId, ibvAddressList[nicIndex]); + if (distance >= 0 && distance < minDistance) { minDistance = distance; - closestIdx = nicIndex; + closestNicIdxs.clear(); + closestNicIdxs.insert(nicIndex); + } else if (distance >= 0 && distance == minDistance) { + closestNicIdxs.insert(nicIndex); } } } } - if (closestIdx != -1) { - topo.closestNicsToGpu[gpuIndex].push_back(closestIdx); - assignedCount[closestIdx]++; - } + for (auto idx : closestNicIdxs) + topo.closestNicsToGpu[gpuIndex].push_back(idx); } // Compute the reverse mapping: closest GPU(s) for each NIC @@ -7544,28 +8412,22 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::set closestGpuIdxs = GetNearestDevicesInTree(ibvDeviceList[nicIndex].busId, gpuAddressList); if (closestGpuIdxs.empty()) { - // Fallback: use bus ID distance + // Fallback: use PCIe domain distance int minDistance = std::numeric_limits::max(); - int closestIdx = -1; - for (int gpuIdx = 0; gpuIdx < numGpus; gpuIdx++) { if (gpuAddressList[gpuIdx].empty()) continue; - - int distance = GetBusIdDistance(ibvDeviceList[nicIndex].busId, gpuAddressList[gpuIdx]); + int distance = GetDomainDistance(ibvDeviceList[nicIndex].busId, gpuAddressList[gpuIdx]); if (distance >= 0 && distance < minDistance) { minDistance = distance; - closestIdx = gpuIdx; + closestGpuIdxs.clear(); + closestGpuIdxs.insert(gpuIdx); + } else if (distance >= 0 && distance == minDistance) { + closestGpuIdxs.insert(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); - } + } + for (int idx : closestGpuIdxs) { + topo.closestGpusToNic[nicIndex].push_back(idx); } } } @@ -7846,7 +8708,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) return {ERR_FATAL, "CPU index must be between 0 and %d inclusively", numCpus - 1}; agent = cpuAgents[exeDevice.exeIndex]; break; - case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: + case EXE_GPU_GFX: case EXE_GPU_DMA: case EXE_GPU_BDMA: case EXE_GPU_TDM: if (exeIndex < 0 || exeIndex >= numGpus) return {ERR_FATAL, "GPU index must be between 0 and %d inclusively", numGpus - 1}; agent = gpuAgents[exeIndex]; @@ -7917,12 +8779,14 @@ const auto& AmdSmiFabricInfoV1(const T& info) std::map cpuAgentMap; hsa_iterate_agents(cpuAgentCallback, &cpuAgentMap); + // Index CPU agents by logical CPU index (physical node = cpuNumaMap[logical]) cpuAgents.clear(); - int numCpus = numa_num_configured_nodes(); + int numCpus = static_cast(cpuNumaMap.size()); cpuAgents.resize(numCpus); for (int i = 0; i < numCpus; i++) { - if (cpuAgentMap.count(i)) { - cpuAgents[i] = cpuAgentMap[i]; + int physNode = cpuNumaMap[i]; + if (cpuAgentMap.count(physNode)) { + cpuAgents[i] = cpuAgentMap[physNode]; } } } @@ -8120,6 +8984,11 @@ const auto& AmdSmiFabricInfoV1(const T& info) return System::Get().GetClosestCpuNumaToGpu(gpuIndex, targetRank); } + int GetCpuNumaPhysicalNode(int cpuIndex) + { + return System::Get().GetCpuPhysicalNode(cpuIndex); + } + int GetClosestCpuNumaToNic(int nicIndex, int targetRank) { return System::Get().GetClosestCpuNumaToNic(nicIndex, targetRank); @@ -8222,6 +9091,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) // Enumerations #undef hipDeviceAttributeClockRate #undef hipDeviceAttributeMultiprocessorCount +#undef hipDeviceAttributeMaxSharedMemoryPerBlock #undef hipDeviceAttributeWarpSize #undef hipErrorPeerAccessAlreadyEnabled #undef hipFuncCachePreferShared @@ -8252,6 +9122,7 @@ const auto& AmdSmiFabricInfoV1(const T& info) #undef hipGetDeviceCount #undef hipGetDeviceProperties #undef hipGetErrorString +#undef hipGetLastError #undef hipHostFree #undef hipHostMalloc #undef hipMalloc @@ -8276,10 +9147,6 @@ const auto& AmdSmiFabricInfoV1(const T& info) #undef hipMemImportFromShareableHandle #endif -// Kernel macros -#undef GetHwId -//#undef GetXccId - // Undefine helper macros #undef ERR_CHECK #undef ERR_APPEND diff --git a/src/header/tdmCopy.h b/src/header/tdmCopy.h new file mode 100644 index 00000000..bdaef85c --- /dev/null +++ b/src/header/tdmCopy.h @@ -0,0 +1,718 @@ +/* +Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +/// \file tdmCopy.h +/// \brief Helper functions to perform HBM to HBM copies via the Tensor Data Mover +/// which copies HBM to a LDS staging buffer, then from LDS out to HBM +/// all without touching cache +/// +/// This is introduced via __device__ level memcpy-like API which can be +/// either blocking or asynchronous, and utilize all warps, or a team of +/// contiguous warps, to allow for other warps to do other tasks +/// +/// \par Quick start +/// \code +/// #include "tdmCopy.h" +/// __shared__ uint8_t staging[N]; // or dynamic extern __shared__ +/// tdm::tdmCopy(dst, src, bytes, staging, N); // block-collective, blocking +/// __syncthreads(); // block-wide visibility +/// \endcode +/// +/// \par Two axes of control +/// - Completion: tdmCopy() (blocking) vs. tdmCopyAsync() + tdmWait() (deferred). +/// - Participation: block-collective (all warps) vs. tdmCopyByTeam() (a +/// contiguous warp range, leaving the other warps free to compute). +/// +/// \par Availability +/// TDM is a hardware feature present on some architectures only. On a target +/// without it, every entry point is `= delete`d: including the header is always +/// fine, but calling any tdm:: function is a hard compile-time error at the call +/// site. See the AVAILABILITY block below. For a runtime/host check (e.g. to pick +/// a kernel before launch), use IsTdmCopySupported(), which is always callable. +/// +/// The design rationale, hardware model, and visibility rules live at the bottom +/// of this file under "IMPLEMENTATION NOTES"; the public API is right here. +#pragma once + +// Two backends are supported: +// AMD's Tensor Data Mover (gfx1250) and NVIDIA's cp.async.bulk / TMA (Hopper, sm_90+). +// The active backend is chosen from the compiler in use (see the AVAILABILITY block below). +#if defined(__CUDACC__) || defined(__NVCC__) || defined(__CUDA__) +# include +# include +#else +# include +#endif +#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 + +// ---- backend selection ----------------------------------------------------- +// TDM_PLATFORM_NV is a host-evaluable proxy for "this is the NVIDIA toolchain". +// Exactly one backend is enabled: TDM_BACKEND_AMD (gfx1250 TDM) or +// TDM_BACKEND_NV (Hopper+ cp.async.bulk). TDM_SUPPORTED is their OR. +#if defined(__CUDACC__) || defined(__NVCC__) || defined(__CUDA__) +# define TDM_PLATFORM_NV 1 +#else +# define TDM_PLATFORM_NV 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 !TDM_PLATFORM_NV +// AMD: TDM builds only when the arch, the builtin, AND the D# descriptor header +// are all present, so an older toolchain degrades gracefully (see rationale above). +# if __has_include() +# define TDM_TOOLCHAIN_AVAILABLE 1 +# else +# define TDM_TOOLCHAIN_AVAILABLE 0 +# endif +# if defined(__gfx1250__) && \ + __has_builtin(__builtin_amdgcn_tensor_load_to_lds) && \ + TDM_TOOLCHAIN_AVAILABLE + /* extend: || (defined(__gfxNNNN__) && ...) */ +# define TDM_BACKEND_AMD 1 +# else +# define TDM_BACKEND_AMD 0 +# endif +# define TDM_BACKEND_NV 0 +#else +// NVIDIA: cp.async.bulk (TMA) is present on Hopper and newer. __CUDA_ARCH__ is +// only defined in the device pass, so this is 0 in the host pass (as intended; +// the host-side IsTdmCopySupported() check below uses the runtime instead). +# define TDM_TOOLCHAIN_AVAILABLE 0 +# define TDM_BACKEND_AMD 0 +# if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900) +# define TDM_BACKEND_NV 1 +# else +# define TDM_BACKEND_NV 0 +# endif +#endif + +#define TDM_SUPPORTED (TDM_BACKEND_AMD || TDM_BACKEND_NV) + +#if TDM_BACKEND_AMD +# include // D# descriptor types for the target's TDM +#endif + +#if TDM_SUPPORTED +# 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__) || defined(__CUDA_ARCH__) + (void)deviceId; + return TDM_SUPPORTED; // compile-time constant for this arch pass +#elif TDM_PLATFORM_NV + // Host (NVIDIA): cp.async.bulk (TMA) requires compute capability 9.0+ (Hopper). + cudaDeviceProp prop; + if (cudaGetDeviceProperties(&prop, deviceId) != cudaSuccess) return false; + return prop.major >= 9; +#else + hipDeviceProp_t prop; + if (hipGetDeviceProperties(&prop, deviceId) != hipSuccess) return false; + // gcnArchName looks like "gfx1250:sramecc+:xnack-"; match the arch prefix. + // Keep this list in sync with the TDM_SUPPORTED arch condition above. + const char* arch = prop.gcnArchName; + const char* p = "gfx1250"; + while (*p && *arch == *p) { ++arch; ++p; } + // Require BOTH a TDM-capable arch AND a toolchain that actually built TDM + // (otherwise the device pass emitted a no-op stub and enabling the path here + // would silently produce wrong results). + return TDM_TOOLCHAIN_AVAILABLE && (*p == '\0'); +#endif +} + +/// \brief Blocking block-collective copy of [0, sizeBytes): dst <- src. +/// +/// Call from ALL threads of the block with identical arguments (memcpy-style +/// order). Work is partitioned across every warp in the block. On return, the +/// calling warp's TDM ops are complete. +/// +/// \param dst Destination in global memory (HBM). +/// \param src Source in global memory (HBM). +/// \param sizeBytes Number of bytes to copy. +/// \param ldsBuffer Per-block LDS staging area (shared memory). +/// \param ldsBufferBytes Size of \p ldsBuffer in bytes; subdivided among warps. +/// +/// \note For BLOCK-wide visibility, follow with __syncthreads() (and see the +/// visibility notes at the bottom of the file). +/// \warning \p src and \p dst must be GLOBAL (HBM) pointers. Passing a shared / +/// LDS pointer compiles (it decays to void*) but is undefined at run +/// time -- the address is programmed into the descriptor's global field. +__device__ TDM_API void tdmCopy(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes) TDM_DELETED; + +/// \brief Non-blocking variant of tdmCopy(): issue and return. +/// +/// Identical partitioning to tdmCopy(), but does NOT drain on return. The last +/// few TDM ops (bounded by the per-wave queue depth) stay in flight so they +/// overlap with whatever the calling warp does next. Pair with tdmWait(). +/// +/// \see tdmCopy for parameter meanings. +/// \see tdmWait to complete the copy. +__device__ TDM_API void tdmCopyAsync(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes) TDM_DELETED; + +/// \brief Blocking WARP-SPECIALIZED copy performed by one contiguous warp team. +/// +/// Only warps in the half-open range [\p startWarpId, \p stopWarpId) participate; +/// all other warps return immediately and are free to do compute. Work and LDS +/// are partitioned by RANK WITHIN THE TEAM (warpId - startWarpId), so several +/// disjoint teams can each run a different copy concurrently, each with its own +/// dst/src and its own \p ldsBuffer region. +/// +/// \param dst Destination in global memory (HBM). +/// \param src Source in global memory (HBM). +/// \param sizeBytes Number of bytes to copy. +/// \param ldsBuffer THIS team's LDS region (distinct per team). +/// \param ldsBufferBytes Size of this team's LDS region, split among its warps. +/// \param startWarpId First warp of the team (inclusive). +/// \param stopWarpId One past the last warp (clamped to nWarps; ~0u = end). +/// +/// \note This does NOT synchronize the block -- you choose the barrier. A single +/// __syncthreads() after the copy/compute branches is enough for +/// independent compute; use named/arrive-wait barriers for a pipelined +/// producer/consumer so the copy team can run ahead. +/// \warning Give each team (and the compute warps) a NON-OVERLAPPING LDS region; +/// the library trusts the pointer/size you pass. +/// +/// \code +/// const uint32_t warpId = threadIdx.x / warpSize; // (1D block) +/// if (warpId < COPY_WARPS) +/// tdm::tdmCopyByTeam(dst, src, n, staging, teamLdsBytes, 0u, COPY_WARPS); +/// else +/// /* compute -- use LDS past the team's window */; +/// __syncthreads(); +/// \endcode +__device__ TDM_API void tdmCopyByTeam(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) TDM_DELETED; + +/// \brief Non-blocking variant of tdmCopyByTeam(): issue and return. +/// \see tdmCopyByTeam for parameter meanings and team semantics. +/// \see tdmWait to complete the copy (called by each participating warp). +__device__ TDM_API void tdmCopyAsyncByTeam(void* dst, const void* src, size_t sizeBytes, + void* ldsBuffer, size_t ldsBufferBytes, + uint32_t startWarpId, uint32_t stopWarpId) TDM_DELETED; + +/// \brief Drain the CALLING WARP's outstanding TDM ops (TENSORcnt -> 0). +/// +/// TENSORcnt is a per-wave counter, so this waits only on the ops this warp +/// issued -- nothing else. +/// +/// \note Multiple teams need no special handling: each participating warp calls +/// tdmWait() to drain its own ops, and one team's wait has no effect on +/// another's (there is no shared counter). Every ISSUING warp must call it +/// -- a warp cannot drain its teammates' ops. tdmCopy*()/*ByTeam() blocking +/// forms already do this internally. +/// \note For BLOCK-wide visibility follow with __syncthreads(); if copied data is +/// consumed within the block AND the vector head ran, also +/// __threadfence_block() so those ordinary global stores are observed. +__device__ TDM_API void tdmWait() TDM_DELETED; + +} // namespace tdm + +#undef TDM_API +#undef TDM_DELETED + + +// ############################################################################ +// # # +// # IMPLEMENTATION NOTES # +// # # +// ############################################################################ +// +// LAYOUT OF A COPY +// [ head (VECTOR) ][ ---- aligned 256B rows (TDM) ---- ][ tail (TDM 1-D) ] +// * Fixed choices for bandwidth: 4-byte data_size, 256B TDM row width. +// * `head` brings the SOURCE up to a 128B boundary (the direct-copy +// requirement); being the unaligned remainder, it stays a cooperative +// vector copy done by the team's first warp. +// * The aligned bulk is 2D TDM tiles (64 dwords x N rows of 256B). +// * `tail` (< 256B sub-row remainder) is a separate 1-D TDM op at byte +// granularity. TDM's out-of-bounds clamp is per-dimension (rectangular) and +// cannot express "N full rows + a partial row", so the partial row must be +// its own tile rather than riding the bulk descriptor's clamp. +// +// HARDWARE MODEL (why the per-tile waits are REQUIRED) +// * TDM engines: 1 per SIMD-pair -> 2 per WGP; a warp runs on 1 SIMD32 and +// shares its pair's engine. Bandwidth needs >=2 issuing warps/block (both +// engines) and many blocks (many WGPs); a lone warp is latency-bound. +// * Same-wave TDM ops ISSUE in order but their memory effects OVERLAP: up to 3 +// ops are outstanding per wave (that is what TENSORcnt counts). In-order issue +// does NOT serialize completion, so it does NOT make LDS reuse hazard-free. +// * This copy is single-buffered (one LDS window per warp), so each tile is a +// load->store dependency chain on that window: the store reads what the load +// wrote (RAW), and the next tile's load overwrites the window the store is +// still draining (WAR). Both edges need an s_wait_tensorcnt, so issueRow()/ +// issueBytes() 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 ================ + +#if TDM_BACKEND_AMD // ----- AMD Tensor Data Mover (gfx1250) -------- + +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 aligned bulk as a single 1-D tile through ONE LDS window. -------- +// The chunk is expressed as one flat 1-D tile of `rows * TD0` dwords (dataSize +// DS4), moving exactly `rows * WIDTH` contiguous bytes through the same +// single-buffered LDS window with the same RAW/WAR waits as the tail path. +// tileDim0 is a 16-bit field, so `rows*TD0` must be <= 65535 (rows <= 1023); +// the per-warp LDS window bounds chunkRows well below that. +__device__ inline void issueRow(uint64_t src, uint32_t lds, uint64_t dst, + uint32_t rows) { + uint32_t dwords = rows * TD0; // whole chunk as one flat dword run + gfx1250_TDM_GROUP1 g1; + g1.dataSize(DS4); + g1.tileDim0(dwords); g1.tileDim1(1); + g1.tensorDim0(dwords); g1.tensorDim1(1); + g1.tensorDim0Stride(dwords); + + 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 issueBytes(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 + issueBytes(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; + issueRow(sBase + off, myLds, dBase + off, chunkRows); + } +} + +} // namespace detail + +#elif TDM_BACKEND_NV // ----- NVIDIA cp.async.bulk / TMA (sm_90+) ----- + +// The AMD backend expresses the aligned bulk as a 2D tensor tile (256B rows). +// cp.async.bulk is a FLAT 1-D byte copy, so the 2D descriptor machinery is gone: +// we copy contiguous 16B-aligned byte ranges through the same single-buffered LDS +// staging window. The load/store completion model also differs from AMD's single +// TENSORcnt: the global->shared load is tracked by an mbarrier (transaction bytes) +// and the shared->global store by a bulk async-group (commit + wait). The +// partition/team logic mirrors the AMD path so the public API is identical. +namespace detail { + +namespace ptx = cuda::ptx; + +constexpr uint32_t ALIGN = 16; // cp.async.bulk addr/size granularity + +// ---- cooperative vector copy of a small byte range, by one warp (== AMD) ---- +__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]; +} + +// Drain the CALLING THREAD's outstanding bulk-store groups (the WAR/async edge). +// Named waitTensor0() so the shared public-API wrappers below are backend-agnostic. +__device__ inline void waitTensor0() { + ptx::cp_async_bulk_wait_group_read(ptx::n32_t<0>{}); +} + +// ---- issue one contiguous chunk (nbytes, a 16B multiple) through ONE window -- +// Single-buffered, so the same RAW/WAR edges as the AMD path apply, just with the +// cp.async.bulk completion mechanisms: +// * G2S load -> mbarrier wait (RAW: store must see the filled LDS window) +// * S2G store -> bulk-group wait (WAR: next load must not clobber a window +// whose store is still draining) +// Issued by the warp's leader thread only (bulk copies are single-thread ops). +__device__ inline void issueChunk(const uint8_t* src, void* lds, uint8_t* dst, + uint32_t nbytes, uint64_t* bar, uint32_t& phase) { + // G2S: arrive + expect nbytes, launch the async copy, then wait for the + // mbarrier phase to flip (the copy deposits its tx-bytes on completion). + ptx::mbarrier_arrive_expect_tx(ptx::sem_release, ptx::scope_cta, + ptx::space_shared, bar, nbytes); + ptx::cp_async_bulk(ptx::space_cluster, ptx::space_global, lds, src, nbytes, bar); + while (!ptx::mbarrier_try_wait_parity(bar, phase)) {} + phase ^= 1u; + + // Order the async-proxy LDS writes before the async-proxy store reads them. + ptx::fence_proxy_async(ptx::space_shared); + + // S2G: launch the store, commit the bulk group, and drain before window reuse. + ptx::cp_async_bulk(ptx::space_global, ptx::space_shared, dst, lds, nbytes); + ptx::cp_async_bulk_commit_group(); + ptx::cp_async_bulk_wait_group_read(ptx::n32_t<0>{}); +} + +// ---- core: partition + issue the whole copy for the team [start, stop). ------ +// Same partitioning contract as the AMD issue(): NO final wait beyond the +// per-chunk drains, work + LDS split by rank within the team. +__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 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; + 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; + const uint32_t rank = warpId - startWarpId; + const uint32_t teamWarps = teamStop - startWarpId; + const uint32_t warpThreads = (nThreads - warpId * W < W) ? (nThreads - warpId * W) : W; + const bool leader = (warpThread == 0); + + // Carve per-team mbarriers (one per warp in the team, indexed by rank) from the + // FRONT of this team's ldsBuffer; the staging windows use the remainder. Keeping + // the barriers in the passed-in dynamic LDS avoids any static __shared__, which + // would otherwise push a full-size dynamic allocation past the per-block cap and + // make the launch fail with "invalid argument". + uint8_t* ldsBase = reinterpret_cast(ldsBuffer); + uint32_t barRegion = ((teamWarps * static_cast(sizeof(uint64_t))) + + (ALIGN - 1)) & ~(ALIGN - 1); + uint64_t* bars = reinterpret_cast(ldsBase); + uint8_t* winBase = ldsBase + barRegion; + uint32_t winBytes = (ldsBytes > barRegion) ? (ldsBytes - barRegion) : 0u; + + // --- split the range: [head][ 16B-aligned bulk ][tail] ------------------- + // cp.async.bulk requires BOTH src and dst 16B-aligned. A single head can only + // align both if they share the same 16B phase; if they don't, there is no + // valid bulk split, so fall back to a pure cooperative vector copy. + uint64_t sAddr = reinterpret_cast(s); + uint64_t dAddr = reinterpret_cast(d); + if (((sAddr ^ dAddr) & (ALIGN - 1)) != 0) { + if (rank == 0) warpVecCopy(s, d, sizeBytes, warpThread, warpThreads); + return; + } + uint32_t head = static_cast((ALIGN - (sAddr & (ALIGN - 1))) & (ALIGN - 1)); + if (head > sizeBytes) head = static_cast(sizeBytes); + size_t rest = sizeBytes - head; + size_t bulk = rest & ~static_cast(ALIGN - 1); // whole 16B units + uint32_t tail = static_cast(rest - bulk); // < 16B remainder + size_t tailOff = head + bulk; + + // --- edges (team's FIRST warp = rank 0): vector head and tail ------------ + if (rank == 0 && head) warpVecCopy(s, d, head, warpThread, warpThreads); + if (rank == 0 && tail) warpVecCopy(s + tailOff, d + tailOff, tail, + warpThread, warpThreads); + + // --- aligned bulk via cp.async.bulk -------------------------------------- + if (bulk == 0) return; + uint32_t maxByLds = winBytes / ALIGN; // #warps we can give a window + if (maxByLds == 0) { // no window room: vector fallback + if (rank == 0) + warpVecCopy(s + head, d + head, bulk, warpThread, warpThreads); + return; + } + uint32_t issuers = teamWarps < maxByLds ? teamWarps : maxByLds; + uint32_t window = (winBytes / issuers) & ~(ALIGN - 1); // per-warp 16B-multiple + if (rank >= issuers) return; // this warp doesn't issue + + // distribute the bulk across issuers by team rank (contiguous 16B units) + size_t units = bulk / ALIGN; + size_t base = units / issuers; + size_t extra = units % issuers; + size_t myUnits = base + (rank < extra ? 1u : 0u); + size_t myStart = rank * base + (rank < extra ? rank : extra); + if (myUnits == 0) return; + + size_t myBytes = myUnits * ALIGN; + uint8_t* myLds = winBase + (size_t)rank * window; + const uint8_t* sBase = s + head + myStart * ALIGN; + uint8_t* dBase = d + head + myStart * ALIGN; + + uint64_t* bar = &bars[rank]; + uint32_t phase = 0; + if (leader) ptx::mbarrier_init(bar, 1); // one arrival (the leader) + __syncwarp(); + + if (leader) { + for (size_t off = 0; off < myBytes; off += window) { + uint32_t chunk = (myBytes - off < window) + ? static_cast(myBytes - off) : window; + issueChunk(sBase + off, myLds, dBase + off, chunk, bar, phase); + } + } +} + +} // namespace detail + +#endif // backend selection + +// ---- 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