Refine the get_local_id/... builtins.

As we could prepare correct value on runtime library side and give
a correct value in the payload for dim 0, 1 and 2. So for these 3
dim argument, we don't need to check it whether in the valid range,
we just read the payload's value.

This way, we can avoid any unecessary branching for normal usage of
these builtin functions. And could avoid a known bool related bug.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Tested-by: Sun, Yi <yi.sun@intel.com>

diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
index 3b191ab..afc5b85 100644 (file)
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
@@ -4379,11 +4379,9 @@ DECL_INTERNAL_WORK_ITEM_FN(get_num_groups)
 #define DECL_PUBLIC_WORK_ITEM_FN(NAME, OTHER_RET)    \
 INLINE unsigned NAME(unsigned int dim) {             \
   if (dim == 0) return __gen_ocl_##NAME##0();        \
-  else if (dim > 0 && dim < get_work_dim()) {        \
-    if (dim == 1) return __gen_ocl_##NAME##1();      \
-    else if (dim == 2) return __gen_ocl_##NAME##2(); \
-  }                                                  \
-  return OTHER_RET;                                  \
+  else if (dim == 1) return __gen_ocl_##NAME##1();   \
+  else if (dim == 2) return __gen_ocl_##NAME##2();   \
+  else return OTHER_RET;                             \
 }
 
 DECL_PUBLIC_WORK_ITEM_FN(get_group_id, 0)

diff --git a/src/cl_api.c b/src/cl_api.c
index f7db4bc..bb09c07 100644 (file)
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -1581,7 +1581,7 @@ clEnqueueNDRangeKernel(cl_command_queue  command_queue,
 {
   size_t fixed_global_off[] = {0,0,0};
   size_t fixed_global_sz[] = {1,1,1};
-  size_t fixed_local_sz[] = {16,1,1};
+  size_t fixed_local_sz[] = {1,1,1};
   cl_int err = CL_SUCCESS;
   cl_uint i;