Add a threshold for forcing ROCr to take blit path

This workaround is to avoid performance penalty of SDMA engine
taking a while to clock up from a lower DPM state. Add env var
GPU_FORCE_BLIT_COPY_SIZE (1024 by default for HIP in KB). Forcing
Src and Dst agent to be amdgpu makes ROCr take blit copy path for
what otherwise should have been SDMA copy

Change-Id: I222f687155f86000d17d66d25182e490b6710463


[ROCm/clr commit: 5f64e6e7ad]

projects/clr/rocclr/device/device.cpp

@@ -486,6 +486,10 @@ Settings::Settings() : value_(0) {
     if (flagIsDefault(HIP_HIDDEN_FREE_MEM)) {
       HIP_HIDDEN_FREE_MEM = 320;
     }
+
+    if (flagIsDefault(GPU_FORCE_BLIT_COPY_SIZE)) {
+      GPU_FORCE_BLIT_COPY_SIZE = 1024;
+    }
   }
 }
 

projects/clr/rocclr/device/rocm/rocblit.cpp

@@ -599,6 +599,13 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
     dstAgent = dstMemory.dev().getBackendDevice();
   }
 
+  // This workaround is needed for performance to get around the slowdown
+  // caused to SDMA engine powering down if its not active. Forcing agents
+  // to amdgpu device causes rocr to take blit path internally.
+  if (size[0] <= dev().settings().sdmaCopyThreshold_) {
+    srcAgent = dstAgent = dev().getBackendDevice();
+  }
+
   const hsa_signal_value_t kInitVal = 1;
   hsa_signal_store_relaxed(completion_signal_, kInitVal);
 
@@ -656,13 +663,19 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
   // Allocate requested size of memory
   while (totalSize > 0) {
     size = std::min(totalSize, dev().settings().stagedXferSize_);
-    hsa_signal_store_relaxed(completion_signal_, kInitVal);
+    hsa_signal_silent_store_relaxed(completion_signal_, kInitVal);
 
     // Copy data from Host to Device
     if (hostToDev) {
+      // This workaround is needed for performance to get around the slowdown
+      // caused to SDMA engine powering down if its not active. Forcing agents
+      // to amdgpu device causes rocr to take blit path internally.
+      const hsa_agent_t srcAgent =
+          (size <= dev().settings().sdmaCopyThreshold_) ? dev().getBackendDevice() : dev().getCpuAgent();
+
       memcpy(hsaBuffer, hostSrc + offset, size);
       status = hsa_amd_memory_async_copy(hostDst + offset, dev().getBackendDevice(), hsaBuffer,
-                                         dev().getCpuAgent(), size, 0, nullptr, completion_signal_);
+                                         srcAgent, size, 0, nullptr, completion_signal_);
       if (status == HSA_STATUS_SUCCESS) {
         hsa_signal_value_t val = hsa_signal_wait_acquire(
             completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
@@ -680,9 +693,15 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
       continue;
     }
 
+    // This workaround is needed for performance to get around the slowdown
+    // caused to SDMA engine powering down if its not active. Forcing agents
+    // to amdgpu device causes rocr to take blit path internally.
+    const hsa_agent_t dstAgent =
+        (size <= dev().settings().sdmaCopyThreshold_) ? dev().getBackendDevice() : dev().getCpuAgent();
+
     // Copy data from Device to Host
     status =
-        hsa_amd_memory_async_copy(hsaBuffer, dev().getCpuAgent(), hostSrc + offset,
+        hsa_amd_memory_async_copy(hsaBuffer, dstAgent, hostSrc + offset,
                                   dev().getBackendDevice(), size, 0, nullptr, completion_signal_);
     if (status == HSA_STATUS_SUCCESS) {
       hsa_signal_value_t val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ,

projects/clr/rocclr/device/rocm/rocsettings.cpp

@@ -74,6 +74,8 @@ Settings::Settings() {
   pinnedXferSize_ = std::min(GPU_PINNED_XFER_SIZE, MaxPinnedXferSize) * Mi;
   pinnedMinXferSize_ = std::min(GPU_PINNED_MIN_XFER_SIZE * Ki, pinnedXferSize_);
 
+  sdmaCopyThreshold_ = GPU_FORCE_BLIT_COPY_SIZE * Ki;
+
   // Don't support Denormals for single precision by default
   singleFpDenorm_ = false;
 

projects/clr/rocclr/device/rocm/rocsettings.hpp

@@ -73,6 +73,8 @@ class Settings : public device::Settings {
   size_t pinnedXferSize_;     //!< Pinned buffer size for transfer
   size_t pinnedMinXferSize_;  //!< Minimal buffer size for pinned transfer
 
+  size_t sdmaCopyThreshold_;  //!< Use SDMA to copy above this size
+
   //! Default constructor
   Settings();
 

projects/clr/rocclr/utils/flags.hpp

@@ -72,6 +72,8 @@ release(bool, GPU_FLUSH_ON_EXECUTION, false,                                  \
         "Submit commands to HW on every operation. 0 - Disable, 1 - Enable")  \
 release(bool, GPU_USE_SYNC_OBJECTS, true,                                     \
         "If enabled, use sync objects instead of polling")                    \
+release(size_t, GPU_FORCE_BLIT_COPY_SIZE, 0,                                  \
+        "Size in KB of the threshold below which to force blit instead for sdma") \
 release(bool, CL_KHR_FP64, true,                                              \
         "Enable/Disable support for double precision")                        \
 release(cstring, AMD_OCL_BUILD_OPTIONS, 0,                                    \