diff --git a/src/operators/kernel/cl/elementwise_add_kernel.cpp b/src/operators/kernel/cl/elementwise_add_kernel.cpp
index a2a20ec09075df6c197a2f7abf0ab70ae111e7d7..0bc3700350f5d54a16fff6b1fc297695d94b769b 100644
--- a/src/operators/kernel/cl/elementwise_add_kernel.cpp
+++ b/src/operators/kernel/cl/elementwise_add_kernel.cpp
@@ -78,9 +78,11 @@ void ElementwiseAddKernel<GPU_CL, float>::Compute(
     int width = input->ImageWidth();
     int height = input->ImageHeight();
     size_t global_work_size[2] = {width, height};
+    cl_event out_event = param.Out()->GetClEvent();
+    cl_event wait_event = param.InputX()->GetClEvent();
     status =
         clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 2,
-                               NULL, global_work_size, NULL, 0, NULL, NULL);
+                               NULL, global_work_size, NULL, 1,&wait_event, &out_event);
     CL_CHECK_ERRORS(status);
   } else {
     DLOG << "error:bias dims is error";
diff --git a/src/operators/kernel/cl/feed_kernel.cpp b/src/operators/kernel/cl/feed_kernel.cpp
index 0edb1f2bee2dd65861e665f76bdd28ac0a3913c0..e79e0ed15dfae00ee894ad6239ceadad704ba6bb 100644
--- a/src/operators/kernel/cl/feed_kernel.cpp
+++ b/src/operators/kernel/cl/feed_kernel.cpp
@@ -51,8 +51,11 @@ void FeedKernel<GPU_CL, float>::Compute(const FeedParam<GPU_CL> &param) {
   CL_CHECK_ERRORS(status);
 
   size_t global_work_size[2] = {width, height};
+
+    cl_event out_event = param.Out()->GetClEvent();
+
   status = clEnqueueNDRangeKernel(this->cl_helper_.CLCommandQueue(), kernel, 2,
-                                  NULL, global_work_size, NULL, 0, NULL, NULL);
+                                  NULL, global_work_size, NULL, 0, NULL, &out_event);
   CL_CHECK_ERRORS(status);
 }