fix(gfx): copy subExecParam from device before reading timing data (#323)

Author: fcui-amd

src/header/TransferBench.hpp

@@ -1498,6 +1498,22 @@ namespace {
     return "?";
   }
 
+  // Returns whether a memory allocation can be directly read by host code.
+  static ErrResult GetMemHostAccessibility(MemDevice const& memDevice, bool& hostAccessible)
+  {
+    hostAccessible = IsCpuMemType(memDevice.memType) || memDevice.memType == MEM_MANAGED;
+    if (hostAccessible || !IsGpuMemType(memDevice.memType)) return ERR_NONE;
+
+#if defined(__NVCC__)
+    return ERR_NONE;
+#else
+    int isLargeBar = 0;
+    ERR_CHECK(hipDeviceGetAttribute(&isLargeBar, hipDeviceAttributeIsLargeBar, memDevice.memIndex));
+    hostAccessible = (isLargeBar != 0);
+    return ERR_NONE;
+#endif
+  }
+
   // Allocate memory
   static ErrResult AllocateMemory(MemDevice memDevice, size_t numBytes, void** memPtr,
                                   size_t* actualBytes = NULL,
@@ -2943,6 +2959,7 @@ namespace {
 
     // For GPU-Executors
     SubExecParam*              subExecParamGpu;   ///< GPU copy of subExecutor parameters
+    bool                       subExecParamHostAccessible; ///< Host can directly read subExecParamGpu
     vector<hipStream_t>        streams;           ///< HIP streams to launch on
     vector<hipEvent_t>         startEvents;       ///< HIP start timing event
     vector<hipEvent_t>         stopEvents;        ///< HIP stop timing event
@@ -4476,6 +4493,7 @@ namespace {
 #endif
       ERR_CHECK(AllocateMemory({memType, exeDevice.exeIndex}, exeInfo.totalSubExecs * sizeof(SubExecParam),
                                (void**)&exeInfo.subExecParamGpu));
+      ERR_CHECK(GetMemHostAccessibility({memType, exeDevice.exeIndex}, exeInfo.subExecParamHostAccessible));
 
       // Create subexecutor parameter array for entire executor
       exeInfo.subExecParamCpu.clear();
@@ -5435,6 +5453,7 @@ namespace {
                                       int           const  xccDim,
                                       ConfigOptions const& cfg,
                                       int           const  gfxKernelIdx,
+                                      bool          const  subExecParamHostAccessible,
                                       TransferResources&   rss)
   {
     // Determine which kernel to launch
@@ -5491,9 +5510,17 @@ namespace {
         if (cfg.general.recordPerIteration) {
           rss.perIterMsec.push_back(deltaMsec);
           std::set<std::pair<int,int>> CUs;
+          std::vector<SubExecParam> 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(rss.subExecParamGpuPtr[i].xccId,
-                                      GetId(rss.subExecParamGpuPtr[i].hwId)));
+            CUs.insert(std::make_pair(subExecParam[i].xccId,
+                                      GetId(subExecParam[i].hwId)));
           }
           rss.perIterCUs.push_back(CUs);
         }
@@ -5528,6 +5555,7 @@ namespace {
                                                xccDim,
                                                std::cref(cfg),
                                                exeInfo.gfxKernelToUse,
+                                               exeInfo.subExecParamHostAccessible,
                                                std::ref(exeInfo.resources[i])));
       }
       for (auto& asyncTransfer : asyncTransfers)
@@ -5537,7 +5565,8 @@ namespace {
       ExecuteGpuTransfer(iteration, exeInfo.totalSubExecs, exeInfo.subExecParamGpu, exeInfo.streams[0],
                          cfg.gfx.useHipEvents ? exeInfo.startEvents[0] : NULL,
                          cfg.gfx.useHipEvents ? exeInfo.stopEvents[0] : NULL,
-                         xccDim, cfg, exeInfo.gfxKernelToUse, exeInfo.resources[0]);
+                         xccDim, cfg, exeInfo.gfxKernelToUse,
+                         exeInfo.subExecParamHostAccessible, exeInfo.resources[0]);
     }
 
     auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
@@ -5560,6 +5589,15 @@ namespace {
       // 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) {
+        std::vector<SubExecParam> 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<int64_t>::max();
@@ -5568,11 +5606,11 @@ namespace {
 
           for (auto subExecIdx : rss.subExecIdx) {
             if (exeInfo.subExecParamCpu[subExecIdx].N != 0) {
-              minStartCycle = std::min(minStartCycle, exeInfo.subExecParamGpu[subExecIdx].startCycle);
-              maxStopCycle  = std::max(maxStopCycle,  exeInfo.subExecParamGpu[subExecIdx].stopCycle);
+              minStartCycle = std::min(minStartCycle, subExecParam[subExecIdx].startCycle);
+              maxStopCycle  = std::max(maxStopCycle,  subExecParam[subExecIdx].stopCycle);
               if (cfg.general.recordPerIteration) {
-                CUs.insert(std::make_pair(exeInfo.subExecParamGpu[subExecIdx].xccId,
-                                          GetId(exeInfo.subExecParamGpu[subExecIdx].hwId)));
+                CUs.insert(std::make_pair(subExecParam[subExecIdx].xccId,
+                                          GetId(subExecParam[subExecIdx].hwId)));
               }
             }
           }