From 1f1b0621170532ebb22bb33b95bf21232958a269 Mon Sep 17 00:00:00 2001 From: foreman Date: Wed, 13 Jun 2018 11:41:03 -0400 Subject: [PATCH] P4 to Git Change 1567767 by jatang@jatang_rocm_lc2 on 2018/06/13 11:06:09 SWDEV-126897 - Support OpenCL Device Enqueue on ROCm. Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#24 edit --- rocclr/runtime/device/rocm/rocblit.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/rocclr/runtime/device/rocm/rocblit.cpp b/rocclr/runtime/device/rocm/rocblit.cpp index b47407d2e7..ce6561795d 100644 --- a/rocclr/runtime/device/rocm/rocblit.cpp +++ b/rocclr/runtime/device/rocm/rocblit.cpp @@ -2139,7 +2139,10 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, signal1 = schedulerSignal; hsa_signal_store_relaxed(signal1, 1); - sp->scheduler_aql.header = 0; + sp->scheduler_aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); sp->scheduler_aql.setup = 1; sp->scheduler_aql.workgroup_size_x = 1; sp->scheduler_aql.workgroup_size_y = 1;