From dad2209b14ca4c138eacf43b57f6836d42434800 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 27 Apr 2017 19:53:00 -0400 Subject: [PATCH] P4 to Git Change 1403557 by skeely@skeely_HSA_linux on 2017/04/27 19:44:48 SWDEV-119446 - Improve small copy perf in SHOC. Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#18 edit [ROCm/clr commit: 285475fccbd9ad62ba2ad36724a011c6a52708f2] --- .../clr/rocclr/runtime/device/rocm/rocblit.cpp | 18 ++++++++++++++++-- 1 file changed, 16 insertions(+), 2 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp index cebf2e5d5c..f2fe4fef34 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp @@ -516,8 +516,22 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, 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); + hsa_signal_value_t val; + + // Use ACTIVE wait for small transfers. + // Might want to be dependent on also having an idle GPU + // or, if queue is busy, may want to enqueue a blank barrier + // before this and wait BLOCKED on its completion signal, followed + // by ACTIVE on this. + + constexpr size_t small_transfer_size = 4 * Mi; + if (size[0] < small_transfer_size) { + val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, + std::numeric_limits::max(), HSA_WAIT_STATE_ACTIVE); + } else { + val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, + std::numeric_limits::max(), HSA_WAIT_STATE_BLOCKED); + } if (val != (kInitVal - 1)) { LogError("Async copy failed"); status = HSA_STATUS_ERROR;