From 5f64e6e7adc90bf59a77de60d3ec0c2fa2b70502 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Fri, 24 Apr 2020 10:07:50 -0700 Subject: [PATCH] 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 --- rocclr/device/device.cpp | 4 ++++ rocclr/device/rocm/rocblit.cpp | 25 ++++++++++++++++++++++--- rocclr/device/rocm/rocsettings.cpp | 2 ++ rocclr/device/rocm/rocsettings.hpp | 2 ++ rocclr/utils/flags.hpp | 2 ++ 5 files changed, 32 insertions(+), 3 deletions(-) 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, \