From 6edca7aa29e4d3855b63e8871029266c8f88668e Mon Sep 17 00:00:00 2001 From: Mario Kicherer Date: Tue, 7 May 2013 21:57:37 +0200 Subject: [PATCH] clEnqueueNDRangeKernel: fix for segfault caused by empty local_work_size Without this fix, an empty local_work_size that is allowed by specification causes a segfault. Merged the block with a further check below. Signed-off-by: Mario Kicherer Reviewed-by: Zhigang Gong --- src/cl_api.c | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/src/cl_api.c b/src/cl_api.c index c9aba20..1009f9a 100644 --- a/src/cl_api.c +++ b/src/cl_api.c @@ -1073,13 +1073,6 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, goto error; } - /* Local size must be non-null */ - for (i = 0; i < work_dim; ++i) - if (UNLIKELY(local_work_size[i] == 0)) { - err = CL_INVALID_WORK_GROUP_SIZE; - goto error; - } - /* Check offset values. We add a non standard restriction. The offsets must * also be evenly divided by the local sizes */ @@ -1089,16 +1082,16 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, err = CL_INVALID_GLOBAL_OFFSET; goto error; } - if (UNLIKELY(global_work_offset[i] % local_work_size[i])) { + if (UNLIKELY(local_work_size != NULL && global_work_offset[i] % local_work_size[i])) { err = CL_INVALID_GLOBAL_OFFSET; goto error; } } - /* Local sizes must divide global sizes */ + /* Local sizes must be non-null and divide global sizes */ if (local_work_size != NULL) for (i = 0; i < work_dim; ++i) - if (UNLIKELY(global_work_size[i] % local_work_size[i])) { + if (UNLIKELY(local_work_size[i] == 0 || global_work_size[i] % local_work_size[i])) { err = CL_INVALID_WORK_GROUP_SIZE; goto error; } -- 2.7.4