From: Junyan He <junyan...@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...@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
_______________________________________________
Beignet mailing list
Beignet@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/beignet
