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: e2b18eda3e]

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<uint64_t>(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;