[Beignet] [PATCH] Fix two bugs about event.

[junyan.he at intel.com]

From: Junyan He <junyan.he at intel.com>

1. NDrangeKernel need to call cl_event_exec every time.
2. Enqueue Barrier event need to add to queue every time.

Signed-off-by: Junyan He <junyan.he at intel.com>
---
 src/cl_api_event.c  |  1 +
 src/cl_api_kernel.c | 28 ++++++++++++----------------
 2 files changed, 13 insertions(+), 16 deletions(-)

diff --git a/src/cl_api_event.c b/src/cl_api_event.c
index 9207021..5f3a116 100644
--- a/src/cl_api_event.c
+++ b/src/cl_api_event.c
@@ -162,6 +162,7 @@ clEnqueueBarrierWithWaitList(cl_command_queue command_queue,
       err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
       break;
     } else if (e_status == CL_COMPLETE) {
+      cl_command_queue_insert_barrier_event(command_queue, e);
       err = cl_event_exec(e, CL_COMPLETE, CL_FALSE);
       if (err != CL_SUCCESS) {
         break;
diff --git a/src/cl_api_kernel.c b/src/cl_api_kernel.c
index 863b47f..13ea8c0 100644
--- a/src/cl_api_kernel.c
+++ b/src/cl_api_kernel.c
@@ -207,18 +207,16 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
       break;
     }
 
-    int i,j,k;
+    int i, j, k;
     const size_t global_wk_sz_div[3] = {
       fixed_global_sz[0] / fixed_local_sz[0] * fixed_local_sz[0],
       fixed_global_sz[1] / fixed_local_sz[1] * fixed_local_sz[1],
-      fixed_global_sz[2] / fixed_local_sz[2] * fixed_local_sz[2]
-    };
+      fixed_global_sz[2] / fixed_local_sz[2] * fixed_local_sz[2]};
 
     const size_t global_wk_sz_rem[3] = {
       fixed_global_sz[0] % fixed_local_sz[0],
       fixed_global_sz[1] % fixed_local_sz[1],
-      fixed_global_sz[2] % fixed_local_sz[2]
-    };
+      fixed_global_sz[2] % fixed_local_sz[2]};
     cl_uint count;
     count = global_wk_sz_rem[0] ? 2 : 1;
     count *= global_wk_sz_rem[1] ? 2 : 1;
@@ -226,20 +224,18 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
 
     const size_t *global_wk_all[2] = {global_wk_sz_div, global_wk_sz_rem};
     /* Go through the at most 8 cases and euque if there is work items left */
-    for (i = 0; i < 2;i++) {
-      for (j = 0; j < 2;j++) {
+    for (i = 0; i < 2; i++) {
+      for (j = 0; j < 2; j++) {
         for (k = 0; k < 2; k++) {
           size_t global_wk_sz_use[3] = {global_wk_all[k][0], global_wk_all[j][1], global_wk_all[i][2]};
           size_t global_dim_off[3] = {
             k * global_wk_sz_div[0] / fixed_local_sz[0],
             j * global_wk_sz_div[1] / fixed_local_sz[1],
-            i * global_wk_sz_div[2] / fixed_local_sz[2]
-          };
+            i * global_wk_sz_div[2] / fixed_local_sz[2]};
           size_t local_wk_sz_use[3] = {
             k ? global_wk_sz_rem[0] : fixed_local_sz[0],
             j ? global_wk_sz_rem[1] : fixed_local_sz[1],
-            i ? global_wk_sz_rem[2] : fixed_local_sz[2]
-          };
+            i ? global_wk_sz_rem[2] : fixed_local_sz[2]};
           if (local_wk_sz_use[0] == 0 || local_wk_sz_use[1] == 0 || local_wk_sz_use[2] == 0)
             continue;
 
@@ -265,11 +261,11 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue,
           if (event_status < CL_COMPLETE) { // Error happend, cancel.
             err = CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST;
             break;
-          } else if (event_status == CL_COMPLETE) {
-            err = cl_event_exec(e, CL_SUBMITTED, CL_FALSE);
-            if (err != CL_SUCCESS) {
-              break;
-            }
+          }
+
+          err = cl_event_exec(e, (event_status == CL_COMPLETE ? CL_SUBMITTED : CL_QUEUED), CL_FALSE);
+          if (err != CL_SUCCESS) {
+            break;
           }
 
           cl_command_queue_enqueue_event(command_queue, e);
-- 
2.7.4