Async load store alternative path for TDM copy kernel

Makefile

@@ -78,6 +78,18 @@ ifeq ($(filter clean,$(MAKECMDGOALS)),)
     ifneq ($(strip $(ROCM_DEVICE_LIB_PATH)),)
       HIPFLAGS += --rocm-device-lib-path=$(ROCM_DEVICE_LIB_PATH)
     endif
+
+    # gfx1250 TDM descriptor header fallback. src/header/tdmCopy.h keys TDM support
+    # on __has_include(<hip/amd_detail/amd_gfx1250_TDM.h>); some ROCm SDKs omit that
+    # descriptor header, in which case every tdm:: entry point compiles to a deleted
+    # no-op stub and IsTdmCopySupported() reports the GPU as unsupported. When the SDK
+    # lacks the header we add a vendored copy (src/tdm_fallback) to the include search
+    # path. The path is added only when the SDK is missing it, so a real SDK header is
+    # never shadowed.
+    ifeq ($(wildcard $(ROCM_PATH)/include/hip/amd_detail/amd_gfx1250_TDM.h),)
+      $(info - SDK lacks hip/amd_detail/amd_gfx1250_TDM.h; using vendored fallback in ./src/tdm_fallback)
+      CXXFLAGS += -I./src/tdm_fallback
+    endif
   endif
 
   ifeq ($(SINGLE_KERNEL), 1)
@@ -91,6 +103,11 @@ ifeq ($(filter clean,$(MAKECMDGOALS)),)
   endif
   COMMON_FLAGS += -I./src/header -I./src/client -I./src/client/Presets -I./third-party/ibverbs
 
+  # Extra user-supplied preprocessor/compiler flags, e.g.
+  #   make EXTRA_CXXFLAGS=-DTB_DISABLE_HWID_QUERY ...
+  # to stub the hardware-ID reads for functional emulators.
+  COMMON_FLAGS += $(EXTRA_CXXFLAGS)
+
   # libibverbs is loaded dynamically at runtime via dlopen/dlsym (see
   # third-party/ibverbs/IbvDynLoad.hpp), so the build never links against -libverbs
   # and does not require libibverbs-dev to be installed. We only need -ldl so

src/client/Presets/AllToAll.hpp

@@ -55,6 +55,7 @@ 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);
 
   // Check that all ranks have at least the number of GPUs requested
@@ -108,6 +109,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   , "%s TDM executor", useTdmExec ? "Using" : "Not using");
       ev.Print("USE_REMOTE_READ", useRemoteRead, "Using %s as executor", useRemoteRead ? "DST" : "SRC");
       printf("\n");
     }
@@ -121,13 +123,22 @@ int AllToAllPreset(EnvVars&          ev,
     Utils::Print("[ERROR] DMA execution can only be used for copies (A2A_MODE=0)\n");
     return ERR_FATAL;
   }
+  if (useTdmExec && useDmaExec) {
+    Utils::Print("[ERROR] USE_TDM_EXEC and USE_DMA_EXEC are mutually exclusive\n");
+    return ERR_FATAL;
+  }
+  if (useTdmExec && (numSrcs != 1 || numDsts != 1)) {
+    Utils::Print("[ERROR] TDM execution can only be used for copies (A2A_MODE=0)\n");
+    return ERR_FATAL;
+  }
   if (numResults * 2 > numRanks) {
     Utils::Print("[ERROR] Number of extrema results requested exceeds number of ranks.  NUM_RESULTS should be at most half the number of ranks\n");
     return ERR_FATAL;
   }
 
   // Collect the number of GPU devices to use
-  ExeType exeType = useDmaExec ? EXE_GPU_DMA : EXE_GPU_GFX;
+  ExeType exeType = useTdmExec ? EXE_GPU_TDM : (useDmaExec ? EXE_GPU_DMA : EXE_GPU_GFX);
+  char const* exeName = useTdmExec ? "TDM" : (useDmaExec ? "DMA" : "GFX");
 
   std::vector<std::map<std::pair<int, int>, int>> reIndex(numRanks);
   std::vector<Transfer> transfers;
@@ -184,10 +195,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", exeName);
   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, exeName, numSubExecs, numSrcs, numDsts,
                devMemTypeStr.c_str(), numQueuePairs, numRanks);
 
   if (transfers.size() == 0) {

src/client/Presets/AllToAllSweep.hpp

@@ -56,11 +56,20 @@ int AllToAllSweepPreset(EnvVars&          ev,
   int showMinOnly   = EnvVars::GetEnvVar("SHOW_MIN_ONLY",   1);
   int useRemoteRead = EnvVars::GetEnvVar("USE_REMOTE_READ", 0);
   int useSpray      = EnvVars::GetEnvVar("USE_SPRAY",       0);
+  int useTdmExec    = EnvVars::GetEnvVar("USE_TDM_EXEC",    0);
   int verbose       = EnvVars::GetEnvVar("VERBOSE",         0);
 
   std::vector<int> blockList  = EnvVars::GetEnvVarArray("BLOCKSIZES", {256,512,768,1024});
   std::vector<int> unrollList = EnvVars::GetEnvVarArray("UNROLLS", {1,2,3,4,6,8});
   std::vector<int> numSesList = EnvVars::GetEnvVarArray("NUM_SUB_EXECS", {4,8,12,16,24,32});
+  std::vector<int> ldsList    = EnvVars::GetEnvVarArray("LDSBYTES", {0});
+
+  // GFX and TDM expose different "middle" sweep dimensions shown in the second column:
+  //   GFX -> UNROLLS  (cfg.gfx.unrollFactor)
+  //   TDM -> LDSBYTES (cfg.tdm.ldsBytes, per-block LDS staging bytes; 0 = device max)
+  // Each executor lacks the other's knob, so only the relevant list is swept.
+  std::vector<int>& midList  = useTdmExec ? ldsList : unrollList;
+  char const*       midLabel = useTdmExec ? "LDS" : "UnR";
 
   // A2A_MODE may be 0,1,2 or else custom numSrcs:numDsts
   int numSrcs, numDsts;
@@ -94,9 +103,13 @@ int AllToAllSweepPreset(EnvVars&          ev,
     ev.Print("NUM_GPU_DEVICES", numGpus          , "Using %d GPUs", numGpus);
     ev.Print("NUM_SUB_EXECS"  , numSesList.size(), EnvVars::ToStr(numSesList).c_str());
     ev.Print("SHOW_MIN_ONLY"  , showMinOnly      , showMinOnly ? "Showing only slowest GPU results" : "Showing slowest and fastest GPU results");
-    ev.Print("UNROLLS"        , unrollList.size(), EnvVars::ToStr(unrollList).c_str());
+    if (useTdmExec)
+      ev.Print("LDSBYTES"     , ldsList.size()   , "%s (0 = device max)", EnvVars::ToStr(ldsList).c_str());
+    else
+      ev.Print("UNROLLS"      , unrollList.size(), EnvVars::ToStr(unrollList).c_str());
     ev.Print("USE_REMOTE_READ", useRemoteRead    , "Using %s as executor", useRemoteRead ? "DST" : "SRC");
     ev.Print("USE_SPRAY"      , useSpray         , "%s per SubExecutor", useSpray ? "All targets" : "One target");
+    ev.Print("USE_TDM_EXEC"   , useTdmExec       , "Using %s executor", useTdmExec ? "TDM" : "GFX");
     ev.Print("VERBOSE"        , verbose          , verbose ? "Display test results" : "Display summary only");
     printf("\n");
   }
@@ -112,8 +125,18 @@ int AllToAllSweepPreset(EnvVars&          ev,
     exit(1);
   }
 
+  if (useTdmExec && (numSrcs != 1 || numDsts != 1)) {
+    printf("[ERROR] TDM execution can only be used for copies (A2A_MODE=0)\n");
+    exit(1);
+  }
+  if (useTdmExec && useSpray) {
+    printf("[ERROR] Cannot use USE_TDM_EXEC with USE_SPRAY (TDM copies one src/dst pair)\n");
+    exit(1);
+  }
+
   // Collect the number of GPU devices to use
-  ExeType exeType = EXE_GPU_GFX;
+  ExeType exeType = useTdmExec ? EXE_GPU_TDM : EXE_GPU_GFX;
+  char const* exeName = useTdmExec ? "TDM" : "GFX";
 
   std::vector<Transfer> transfers;
 
@@ -180,7 +203,8 @@ int AllToAllSweepPreset(EnvVars&          ev,
     }
   }
 
-  Utils::Print("GPU-GFX All-To-All Sweep benchmark (%lu bytes, local=%s). All values are %s GB/s\n",
+  Utils::Print("GPU-%s All-To-All Sweep benchmark (%lu bytes, local=%s). All values are %s GB/s\n",
+               exeName,
                numBytesPerTransfer,
                a2aLocal         ? "yes"                            : "no",
                ev.useHipEvents  ? "GPU-Event-Timed (min over GPUs)": "CPU-Timed");
@@ -196,10 +220,10 @@ int AllToAllSweepPreset(EnvVars&          ev,
   char sep = ev.outputToCsv ? ',' : ' ';
 
   double bestMinBw = 0.0;
-  int bestBlock = -1, bestUnroll = -1, bestNumSes = -1;
+  int bestBlock = -1, bestMid = -1, bestNumSes = -1;
 
   // Print header once
-  Utils::Print(" BlkS %c UnR ", sep);
+  Utils::Print(" BlkS %c %3s ", sep, midLabel);
   for (int c : numSesList) {
     Utils::Print("%c  SE %03d", sep, c);
     if (ev.useHipEvents && !showMinOnly) {
@@ -208,15 +232,17 @@ int AllToAllSweepPreset(EnvVars&          ev,
   }
   Utils::Print("\n");
 
-  // Results keyed by (blockSize, numSes, unroll) for verbose output
+  // Results keyed by (blockSize, numSes, midValue) for verbose output
   std::map<std::tuple<int,int,int>, TransferBench::TestResults> results;
 
   for (int blockSize : blockList) {
-    cfg.gfx.blockSize = blockSize;
+    if (useTdmExec) cfg.tdm.blockSize = blockSize;
+    else            cfg.gfx.blockSize = blockSize;
 
-    for (int u : unrollList) {
-      cfg.gfx.unrollFactor = u;
-      Utils::Print("%5d %c %3d ", blockSize, sep, u);
+    for (int mid : midList) {
+      if (useTdmExec) cfg.tdm.ldsBytes    = mid;   // per-block LDS staging bytes
+      else            cfg.gfx.unrollFactor = mid;   // GFX-kernel unroll factor
+      Utils::Print("%5d %c %3d ", blockSize, sep, mid);
       fflush(stdout);
 
       for (int c : numSesList) {
@@ -247,11 +273,11 @@ int AllToAllSweepPreset(EnvVars&          ev,
           if (minBw > bestMinBw) {
             bestMinBw  = minBw;
             bestBlock  = blockSize;
-            bestUnroll = u;
+            bestMid    = mid;
             bestNumSes = c;
           }
           if (verbose) {
-            results[std::make_tuple(blockSize, c, u)] = result;
+            results[std::make_tuple(blockSize, c, mid)] = result;
           }
         }
         Utils::Print("%c%8.2f", sep, minBw);
@@ -270,13 +296,13 @@ int AllToAllSweepPreset(EnvVars&          ev,
     int testNum = 0;
     for (int blockSize : blockList) {
       for (int c : numSesList) {
-        for (int u : unrollList) {
+        for (int mid : midList) {
           auto verboseTransfers = transfers;
           for (auto& t : verboseTransfers) {
             t.numSubExecs = useSpray ? (c * targetCount) : c;
           }
-          Utils::Print("BlockSize: %d SubExecs: %d Unroll: %d\n", blockSize, c, u);
-          Utils::PrintResults(ev, ++testNum, verboseTransfers, results[std::make_tuple(blockSize, c, u)]);
+          Utils::Print("BlockSize: %d SubExecs: %d %s: %d\n", blockSize, c, midLabel, mid);
+          Utils::PrintResults(ev, ++testNum, verboseTransfers, results[std::make_tuple(blockSize, c, mid)]);
         }
       }
     }
@@ -287,7 +313,7 @@ int AllToAllSweepPreset(EnvVars&          ev,
     Utils::Print("Highest %s bandwidth found: %7.2f GB/s\n",
                  ev.useHipEvents ? "GPU-event-timed (min)" : "CPU-timed", bestMinBw);
     Utils::Print("          BlockSize  : %7d\n", bestBlock);
-    Utils::Print("          Unroll     : %7d\n", bestUnroll);
+    Utils::Print("          %-10s : %7d\n", midLabel, bestMid);
     Utils::Print("          NumSubExec : %7d\n", bestNumSes);
   }
 

src/header/TransferBench.hpp

@@ -80,6 +80,7 @@ THE SOFTWARE.
 #endif
 
 #include "tdmCopy.h"
+#include "asyncCopy.h"
 
 #endif
 /// @endcond
@@ -713,8 +714,15 @@ namespace TransferBench
 // Helper macro functions
 //==========================================================================================
 
+// TB_DISABLE_HWID_QUERY: stub the hardware-ID reads (CU/SM and XCC) to 0. Some
+// targets (e.g. functional emulators) cannot decode the s_getreg/s_sendmsg_rtn
+// instructions these macros emit; define this at build time to run the kernels
+// there. Real-hardware profiling is unaffected unless this is defined.
+
 // Macro for collecting CU/SM GFX kernel is running on
-#if defined(__GFX9__)
+#if defined(TB_DISABLE_HWID_QUERY)
+  #define GetHwId(hwId) hwId = 0
+#elif defined(__GFX9__)
   #define GetHwId(hwId) asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID)" : "=s" (hwId))
 #elif defined(__GFX10__) || defined(__GFX11__) || defined(__GFX12__)
   #define GetHwId(hwId) asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_HW_ID1)" : "=s" (hwId))
@@ -725,7 +733,9 @@ namespace TransferBench
 #endif
 
 // Macro for collecting XCC GFX kernel is running on
-#if defined(__gfx942__) || defined(__gfx950__)
+#if defined(TB_DISABLE_HWID_QUERY)
+#define GetXccId(val) val = 0
+#elif 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) \
@@ -5850,6 +5860,14 @@ namespace {
 
 // TDM Executor-related functions
 //========================================================================================
+// Compile-time selection of the LDS-staged copy backend used by GpuTdmKernel:
+//   TB_USE_ASYNC_COPY=1 (default) -> async-to/from-LDS path (async::tdmCopy, asyncCopy.h)
+//   TB_USE_ASYNC_COPY=0           -> tensor-data-mover path  (tdm::tdmCopy,  tdmCopy.h)
+// Override without touching source, e.g.: make ... EXTRA_CXXFLAGS=-DTB_USE_ASYNC_COPY=0
+#ifndef TB_USE_ASYNC_COPY
+#define TB_USE_ASYNC_COPY 1
+#endif
+
 #if TDM_SUPPORTED
   __global__ void  GpuTdmKernel(SubExecParam* params,
                                 uint32_t      ldsBytes,
@@ -5869,11 +5887,14 @@ namespace {
     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);
+    constexpr bool useTDM = (TB_USE_ASYNC_COPY == 0);
+    for (int subIterations = 0; subIterations < numSubIterations; subIterations++) {
+      if constexpr (useTDM) {
+        tdm::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes);
+      } else {
+        async::tdmCopy(dst, src, sizeBytes, shmem, ldsBytes);
+      }
       __syncthreads(); // Wait for all warps to finish
-      if (++subIterations == numSubIterations) break;
     }
 
     if (shouldRecordTiming) {