TransferBench v1.69

CHANGELOG.md

@@ -3,6 +3,14 @@
 Documentation for TransferBench is available at
 [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench).
 
+## v1.69.00
+### Added
+- Added support for `ALWAYS_VALIDATE=-1` to disable validation
+### Fixed
+- Fix for non-zero byte offsets used with DMA executor
+- Explicit copy back to host for GPU subExecParam on platforms without large BAR
+- Add missing device synchronization after `hipMemcpy` in `RunTransfers()`
+
 ## v1.68.00
 ### Fixed
 - Improper draining of writes that could artificially inflate transfer timing

src/client/EnvVars.hpp

@@ -79,7 +79,7 @@ class EnvVars
   int useInteractive;                // Pause for user-input before starting transfer loop
 
   // Data options
-  int alwaysValidate;                // Validate after each iteration instead of once after all iterations
+  int alwaysValidate;                // <0 = disable validation, 0 = validate once after all iterations, >0 = validate after each iteration
   int blockBytes;                    // Each subexecutor, except the last, gets a multiple of this many bytes to copy
   int byteOffset;                    // Byte-offset for memory allocations
   vector<float> fillPattern;         // Pattern of floats used to fill source data
@@ -149,7 +149,7 @@ class EnvVars
     else if (archName == "gfx942") defaultGfxUnroll = 4;
     else if (archName == "gfx950") defaultGfxUnroll = 4;
 
-    alwaysValidate    = GetEnvVar("ALWAYS_VALIDATE"     , 0);
+    alwaysValidate    = GetEnvVar("ALWAYS_VALIDATE", 0);
     blockBytes        = GetEnvVar("BLOCK_BYTES"         , 256);
     byteOffset        = GetEnvVar("BYTE_OFFSET"         , 0);
     fillCompress      = GetEnvVarArray("FILL_COMPRESS"  , {});
@@ -352,7 +352,7 @@ class EnvVars
   {
     printf("Environment variables (client):\n");
     printf("======================\n");
-    printf(" ALWAYS_VALIDATE     - Validate after each iteration instead of once after all iterations\n");
+    printf(" ALWAYS_VALIDATE     - Data validation mode: <0=disabled, 0=validate once after all iterations (default), >0=validate after each iteration\n");
     printf(" BLOCK_BYTES         - Controls granularity of how work is divided across subExecutors\n");
     printf(" BYTE_OFFSET         - Initial byte-offset for memory allocations.  Must be multiple of 4\n");
     printf(" CU_MASK             - CU mask for streams. Can specify ranges e.g '5,10-12,14'\n");
@@ -454,7 +454,8 @@ class EnvVars
     if (hideEnv) return;
 
     Print("ALWAYS_VALIDATE", alwaysValidate,
-          "Validating after %s", (alwaysValidate ? "each iteration" : "all iterations"));
+          "Validation %s", (alwaysValidate < 0 ? "disabled" :
+                            alwaysValidate > 0 ? "after each iteration" : "after all iterations"));
     Print("BLOCK_BYTES", blockBytes,
           "Each CU gets a mulitple of %d bytes to copy", blockBytes);
     Print("BYTE_OFFSET", byteOffset,

src/client/Utilities.hpp

@@ -551,6 +551,8 @@ namespace TransferBench::Utils
     if (!RankDoesOutput()) return;
 
     if (!ev.outputToCsv) printf("Test %d:\n", testNum);
+    if (ev.alwaysValidate < 0 && !ev.outputToCsv && !ev.useInteractive)
+      printf("Validation disabled (ALWAYS_VALIDATE < 0)\n");
 
     bool isMultiRank = TransferBench::GetNumRanks() > 1;
 

src/header/TransferBench.hpp

@@ -92,7 +92,7 @@ namespace TransferBench
   using std::set;
   using std::vector;
 
-  constexpr char VERSION[] = "1.68";
+  constexpr char VERSION[] = "1.69";
 
   /**
    * Enumeration of supported Executor types
@@ -217,7 +217,7 @@ namespace TransferBench
    */
   struct DataOptions
   {
-    int           alwaysValidate   = 0;         ///< Validate after each iteration instead of once at end
+    int           alwaysValidate   = 0;         ///< <0 = disable validation, 0 = validate once at end, >0 = validate after each iteration
     int           blockBytes       = 256;       ///< Each subexecutor works on a multiple of this many bytes
     int           byteOffset       = 0;         ///< Byte-offset for memory allocations
     vector<float> fillPattern      = {};        ///< Pattern of floats used to fill source data
@@ -5545,6 +5545,8 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
     int numDsts = (int)resources.dstMem.size();
     ERR_CHECK(hipSetDevice(exeIndex));
     int subIterations = 0;
+    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)
         ERR_CHECK(hipEventRecord(startEvent, stream));
@@ -5558,12 +5560,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
       do {
         // Queue for each output location
         for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) {
+          float* const dst = resources.dstMem[dstIdx] + initOffset;
 #if defined(CUMEM_ENABLED)
-          ERR_CHECK(cuMemcpyAsync((CUdeviceptr)resources.dstMem[dstIdx],
-                                  (CUdeviceptr)resources.srcMem[0],
+          ERR_CHECK(cuMemcpyAsync((CUdeviceptr)dst,
+                                  (CUdeviceptr)src,
                                   resources.numBytes, stream));
 #else
-          ERR_CHECK(hipMemcpyAsync(resources.dstMem[dstIdx], resources.srcMem[0], resources.numBytes,
+          ERR_CHECK(hipMemcpyAsync(dst, src, resources.numBytes,
                                    memcpyKind, stream));
 #endif
         }
@@ -5580,14 +5583,15 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
       do {
         hsa_signal_store_screlease(resources.signal, numDsts);
         for (int dstIdx = 0; dstIdx < numDsts; dstIdx++) {
+          float* const dst = resources.dstMem[dstIdx] + initOffset;
           if (!useSubIndices) {
-            ERR_CHECK(hsa_amd_memory_async_copy(resources.dstMem[dstIdx], resources.dstAgent[dstIdx],
-                                                resources.srcMem[0], resources.srcAgent,
+            ERR_CHECK(hsa_amd_memory_async_copy(dst, resources.dstAgent[dstIdx],
+                                                src, resources.srcAgent,
                                                 resources.numBytes, 0, NULL,
                                                 resources.signal));
           } else {
-            HSA_CALL(hsa_amd_memory_async_copy_on_engine(resources.dstMem[dstIdx], resources.dstAgent[dstIdx],
-                                                         resources.srcMem[0], resources.srcAgent,
+            HSA_CALL(hsa_amd_memory_async_copy_on_engine(dst, resources.dstAgent[dstIdx],
+                                                         src, resources.srcAgent,
                                                          resources.numBytes, 0, NULL,
                                                          resources.signal,
                                                          resources.sdmaEngineId, true));
@@ -5912,24 +5916,30 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
       }
     }
 
-    // Prepare reference src/dst arrays - only once for largest size
+    // Prepare reference src/dst arrays - only once for largest size.
+    // dstReference (expected results) is only needed when validation is enabled.
+    bool const validateEnabled = (cfg.data.alwaysValidate >= 0);
     size_t maxN = maxNumBytes / sizeof(float);
-    vector<float> outputBuffer(maxN);
-    vector<vector<float>> dstReference(maxNumSrcs + 1, vector<float>(maxN));
+    vector<float> outputBuffer(validateEnabled ? maxN : 0);
+    vector<vector<float>> dstReference;
     {
       size_t initOffset = cfg.data.byteOffset / sizeof(float);
       vector<vector<float>> srcReference(maxNumSrcs, vector<float>(maxN));
-      memset(dstReference[0].data(), MEMSET_CHAR, maxNumBytes);
 
+      if (validateEnabled) {
+        dstReference.assign(maxNumSrcs + 1, vector<float>(maxN));
+        memset(dstReference[0].data(), MEMSET_CHAR, maxNumBytes);
+      }
       for (int numSrcs = 0; numSrcs < maxNumSrcs; numSrcs++) {
         PrepareReference(cfg, srcReference[numSrcs], numSrcs);
-        for (int i = 0; i < maxN; i++) {
-          dstReference[numSrcs+1][i] = (numSrcs == 0 ? 0 : dstReference[numSrcs][i]) + srcReference[numSrcs][i];
-        }
+        if (validateEnabled)
+          for (int i = 0; i < maxN; i++)
+            dstReference[numSrcs+1][i] = (numSrcs == 0 ? 0 : dstReference[numSrcs][i]) + srcReference[numSrcs][i];
       }
       // Release un-used partial sums
-      for (int numSrcs = 0; numSrcs < minNumSrcs; numSrcs++)
-        dstReference[numSrcs].clear();
+      if (validateEnabled)
+        for (int numSrcs = 0; numSrcs < minNumSrcs; numSrcs++)
+          dstReference[numSrcs].clear();
 
       // Initialize all src memory buffers (if on local rank)
       bool const verbose = System::Get().IsVerbose();
@@ -6047,7 +6057,7 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
       auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
       double deltaSec = std::chrono::duration_cast<std::chrono::duration<double>>(cpuDelta).count() / cfg.general.numSubIterations;
 
-      if (cfg.data.alwaysValidate) {
+      if (cfg.data.alwaysValidate > 0) {
         ERR_APPEND(ValidateAllTransfers(cfg, transfers, transferResources, dstReference, outputBuffer),
                    errResults);
       }
@@ -6058,11 +6068,32 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
       }
     }
 
-    // Pause for interactive mode - run validation and show per-transfer result before prompt
-    if (cfg.general.useInteractive) {
+    // Interactive per-transfer PASS/FAIL display (skipped when validation disabled). This does its
+    // own comparison pass to show results; authoritative error reporting still comes from
+    // ValidateAllTransfers (mode 0) or the inline per-iteration validation (mode >0).
+    if (cfg.general.useInteractive && cfg.data.alwaysValidate >= 0) {
+      bool inputOk = true;
       if (localRank == 0) {
-        System::Get().Log("Transfers complete. Validation results:\n");
+        if (cfg.data.alwaysValidate == 0) {
+          System::Get().Log("Transfers complete. Hit <Enter> to run validation: ");
+          fflush(stdout);
+          if (getchar() == EOF) {
+            System::Get().Log("[ERROR] Unexpected EOF while waiting for input\n");
+            inputOk = false;
+          } else {
+            System::Get().Log("\nValidation results:\n");
+          }
+        } else {
+          System::Get().Log("Transfers complete.\nValidation results:\n");
+        }
+      }
+
+      // Coordinate any EOF abort so no rank is left waiting at the Barrier below
+      System::Get().Broadcast(0, sizeof(inputOk), &inputOk);
+      if (!inputOk)
+        ERR_APPEND((ErrResult{ERR_FATAL, "Unexpected EOF while waiting for interactive input"}), errResults);
 
+      if (localRank == 0) {
         size_t initOffset = cfg.data.byteOffset / sizeof(float);
         int numPass = 0, numFail = 0;
         for (auto rss : transferResources) {
@@ -6109,20 +6140,17 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
           if (anyLocalDst) { if (transferOk) ++numPass; else ++numFail; }
         }
         System::Get().Log("  Summary: %d PASS  %d FAIL\n", numPass, numFail);
-        System::Get().Log("Hit <Enter> to continue: ");
-        fflush(stdout);
-        if (scanf("%*c") != 0)  {
-          System::Get().Log("[ERROR] Unexpected input\n");
-          exit(1);
-        }
-        System::Get().Log("\n");
         fflush(stdout);
       }
       System::Get().Barrier();
+    } else if (cfg.general.useInteractive) {
+      if (localRank == 0)
+        System::Get().Log("Transfers complete. Validation disabled (ALWAYS_VALIDATE < 0)\n");
+      System::Get().Barrier();
     }
 
     // Validate results
-    if (!cfg.data.alwaysValidate) {
+    if (cfg.data.alwaysValidate == 0) {
       ERR_APPEND(ValidateAllTransfers(cfg, transfers, transferResources, dstReference, outputBuffer),
                  errResults);
     }