From 1bb55cbb6f1b12c13664ecfe695b4a2b5659ae5e Mon Sep 17 00:00:00 2001
From: zhaojiaying01 <zhaojiaying01@baidu.com>
Date: Fri, 19 Oct 2018 20:47:10 +0800
Subject: [PATCH] add wait event for kernel

---
 src/operators/kernel/cl/batchnorm_kernel.cpp | 5 ++++-
 src/operators/kernel/cl/fetch_kernel.cpp     | 3 ++-
 src/operators/kernel/cl/pool_kernel.cpp      | 4 +++-
 3 files changed, 9 insertions(+), 3 deletions(-)

diff --git a/src/operators/kernel/cl/batchnorm_kernel.cpp b/src/operators/kernel/cl/batchnorm_kernel.cpp
index b3ef202782..2da7063796 100644
--- a/src/operators/kernel/cl/batchnorm_kernel.cpp
+++ b/src/operators/kernel/cl/batchnorm_kernel.cpp
@@ -86,8 +86,11 @@ void BatchNormKernel<GPU_CL, float>::Compute(
   clSetKernelArg(kernel, 4, sizeof(cl_mem), &new_bias);
   clSetKernelArg(kernel, 5, sizeof(cl_mem), &out);
 
+  cl_event out_event = param.OutputY()->GetClEvent();
+  cl_event wait_event = param.InputX()->GetClEvent();
   clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
-                         default_work_size.data(), NULL, 0, NULL, NULL);
+                         default_work_size.data(), NULL, 1, &wait_event,
+                         &out_event);
 }
 
 template class BatchNormKernel<GPU_CL, float>;
diff --git a/src/operators/kernel/cl/fetch_kernel.cpp b/src/operators/kernel/cl/fetch_kernel.cpp
index 23c5e8de46..c6f8e78361 100644
--- a/src/operators/kernel/cl/fetch_kernel.cpp
+++ b/src/operators/kernel/cl/fetch_kernel.cpp
@@ -73,8 +73,9 @@ void FetchKernel<GPU_CL, float>::Compute(const FetchParam<GPU_CL> &param) {
     clSetKernelArg(kernel, 6, sizeof(int), &size_batch);
   }
 
+  cl_event wait_event = param.InputX()->GetClEvent();
   clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
-                         default_work_size.data(), NULL, 0, NULL, NULL);
+                         default_work_size.data(), NULL, 1, &wait_event, NULL);
 
   memcpy(out->data<float>(), out_cl_tensor.Data<float>(), out->memory_size());
 }
diff --git a/src/operators/kernel/cl/pool_kernel.cpp b/src/operators/kernel/cl/pool_kernel.cpp
index 802de26e61..904bb0336a 100644
--- a/src/operators/kernel/cl/pool_kernel.cpp
+++ b/src/operators/kernel/cl/pool_kernel.cpp
@@ -63,8 +63,10 @@ void PoolKernel<GPU_CL, float>::Compute(const PoolParam<GPU_CL> &param) {
   clSetKernelArg(kernel, 10, sizeof(cl_mem), &input);
   clSetKernelArg(kernel, 11, sizeof(cl_mem), &out);
 
+  cl_event out_event = param.Output()->GetClEvent();
+  cl_event wait_event = param.Input()->GetClEvent();
   clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 3, NULL,
-                         default_work_size.data(), NULL, 0, NULL, NULL);
+                         default_work_size.data(), NULL, 1, &wait_event, &out_event);
 }
 
 template class PoolKernel<GPU_CL, float>;
-- 
GitLab