From efb43fbae7f9b11e7942e4f563ccd0ca5e2c7d57 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 5 Jul 2018 09:39:08 -0400 Subject: [PATCH] P4 to Git Change 1576542 by jatang@jatang-opencl-hsa-stg5 on 2018/07/05 09:29:34 SWDEV-126897 - Remove un-necessary synchronize() for Device Enqueue in LC/ROCm path. We don't need it since we have a hsa_signal_wait_acquire() to wait for the scheduler to finish. Moreover, synchronize() sends a barrier packet that could get stuck for some reason. Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#25 edit [ROCm/clr commit: e2b18eda3ec9748d7a1f2b289d1df03d18a83507] --- projects/clr/rocclr/runtime/device/rocm/rocblit.cpp | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp index ce6561795d..6af9c14434 100644 --- a/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/runtime/device/rocm/rocblit.cpp @@ -2135,9 +2135,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, sp->child_queue = reinterpret_cast(schedulerQueue); sp->complete_signal = schedulerSignal; - hsa_signal_t signal1; - signal1 = schedulerSignal; - hsa_signal_store_relaxed(signal1, 1); + hsa_signal_store_relaxed(schedulerSignal, 1); sp->scheduler_aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | @@ -2166,9 +2164,8 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, bool result = false; result = gpu().submitKernelInternal(ndrange, *kernels_[Scheduler], parameters, nullptr); releaseArguments(parameters); - synchronize(); - if (hsa_signal_wait_acquire(signal1, HSA_SIGNAL_CONDITION_LT, 1, (-1), + if (hsa_signal_wait_acquire(schedulerSignal, HSA_SIGNAL_CONDITION_LT, 1, (-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogWarning("Failed schedulerSignal wait"); return false;