Fixing GFX Executor missing threadfence_system from v1.67

CHANGELOG.md

@@ -3,6 +3,12 @@
 Documentation for TransferBench is available at
 [https://rocm.docs.amd.com/projects/TransferBench](https://rocm.docs.amd.com/projects/TransferBench).
 
+## v1.68.00
+### Fixed
+- Improper draining of writes that could artificially inflate transfer timing
+- Potential timing bug when running GFX Executor in warp-subexecutor mode with small data sizes
+- Keeping subiteration threads in sync per subiteration
+
 ## v1.67.00
 ### Added
 - Added NIC_TRAFFIC_CLASS to set the DSCP/traffic class byte in the RoCE GRH for QPs (RoCE only)

src/header/TransferBench.hpp

@@ -92,7 +92,7 @@ namespace TransferBench
   using std::set;
   using std::vector;
 
-  constexpr char VERSION[] = "1.67";
+  constexpr char VERSION[] = "1.68";
 
   /**
    * Enumeration of supported Executor types
@@ -5036,21 +5036,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
           }
         }
       }
-      // Allows for numSubiterations == 0 to run infinitely
-      if (++subIterations == numSubIterations) break;
-    }
 
-    // Wait for all threads to finish
-    if (seType == 1) {
-      // For warp-level, sync within warp only
+      // Drain every thread's writes out to system scope
+      __threadfence_system();
+
+      // Wait for all threads to finish this subiteration
+      if (seType == 1) {
+        // For warp-level, sync within warp only
 #if defined(__HIP_PLATFORM_AMD__) && (HIP_VERSION_MAJOR < 7)
-      __builtin_amdgcn_wave_barrier();
+        __builtin_amdgcn_wave_barrier();
 #else
-      __syncwarp();
+        __syncwarp();
 #endif
-    } else {
-      // For threadblock-level, sync all threads
-      __syncthreads();
+      } else {
+        // For threadblock-level, sync all threads
+        __syncthreads();
+      }
+
+      // Allows for numSubiterations == 0 to run infinitely
+      if (++subIterations == numSubIterations) break;
     }
 
     if (shouldRecordTiming) {
@@ -5213,22 +5217,25 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
           }
         }
       }
-      // Allows for numSubiterations == 0 to run infinitely
-      if (++subIterations == numSubIterations) break;
-    }
 
-    // Wait for all threads to finish
-    if (seType == 1) {
-      // For warp-level, sync within warp only
+      // Drain every thread's writes out to system scope
+      __threadfence_system();
+
+      // Wait for all threads to finish this subiteration
+      if (seType == 1) {
+        // For warp-level, sync within warp only
 #if defined(__HIP_PLATFORM_AMD__) && (HIP_VERSION_MAJOR < 7)
-      __builtin_amdgcn_wave_barrier();
+        __builtin_amdgcn_wave_barrier();
 #else
-
-      __syncwarp();
+        __syncwarp();
 #endif
-    } else {
-      // For threadblock-level, sync all threads
-      __syncthreads();
+      } else {
+        // For threadblock-level, sync all threads
+        __syncthreads();
+      }
+
+      // Allows for numSubiterations == 0 to run infinitely
+      if (++subIterations == numSubIterations) break;
     }
 
     if (shouldRecordTiming) {
@@ -5459,11 +5466,13 @@ static bool IsConfiguredGid(union ibv_gid const& gid)
           std::set<std::pair<int, int>> CUs;
 
           for (auto subExecIdx : rss.subExecIdx) {
-            minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle);
-            maxStopCycle  = std::max(maxStopCycle,  exeInfo.subExecParamGpu[subExecIdx].stopCycle);
-            if (cfg.general.recordPerIteration) {
-              CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId,
-                                        GetId(exeInfo.subExecParamGpu[subExecIdx].hwId)));
+            if (exeInfo.subExecParamCpu[subExecIdx].N != 0) {
+              minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle);
+              maxStopCycle  = std::max(maxStopCycle,  exeInfo.subExecParamGpu[subExecIdx].stopCycle);
+              if (cfg.general.recordPerIteration) {
+                CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId,
+                                          GetId(exeInfo.subExecParamGpu[subExecIdx].hwId)));
+              }
             }
           }