From cdb8cc2077441b15bc9a7ec7d77fc12e560de475 Mon Sep 17 00:00:00 2001 From: Zhigang Gong Date: Tue, 25 Jun 2013 14:15:09 +0800 Subject: [PATCH] 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 Tested-by: Sun, Yi --- backend/src/ocl_stdlib.h | 8 +++----- src/cl_api.c | 2 +- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h index 3b191ab..afc5b85 100644 --- 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 --- 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; -- 2.7.4