diff --git a/rocclr/device/device.cpp b/rocclr/device/device.cpp
index 9f7771798e..05be39c8f2 100755
--- a/rocclr/device/device.cpp
+++ b/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;
+    }
   }
 }
 
diff --git a/rocclr/device/rocm/rocblit.cpp b/rocclr/device/rocm/rocblit.cpp
index d5dc5e3a81..6f8be50f37 100755
--- a/rocclr/device/rocm/rocblit.cpp
+++ b/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,
diff --git a/rocclr/device/rocm/rocsettings.cpp b/rocclr/device/rocm/rocsettings.cpp
index ff2606e3f7..62176fe352 100644
--- a/rocclr/device/rocm/rocsettings.cpp
+++ b/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;
 
diff --git a/rocclr/device/rocm/rocsettings.hpp b/rocclr/device/rocm/rocsettings.hpp
index 343aeb2095..8744d2ccb0 100644
--- a/rocclr/device/rocm/rocsettings.hpp
+++ b/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();
 
diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp
index ee2dc77c26..f67aa46e48 100644
--- a/rocclr/utils/flags.hpp
+++ b/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,                                    \