diff options
author | Luo Xionghu <xionghu.luo@intel.com> | 2014-09-10 11:31:32 +0800 |
---|---|---|
committer | Zhigang Gong <zhigang.gong@intel.com> | 2014-09-11 13:42:31 +0800 |
commit | 55ac17eba32cb440f09a0cef63d0675ce9ebbfbb (patch) | |
tree | 60589684e8786a6abfe04d8c461a66da1188cc96 | |
parent | ee4e915ada224fda3fc589da141778be1f9e07ab (diff) |
fix clGetKernelWorkGroupInfo built-in kernel fail.
add CL_KERNEL_GLOBAL_WORK_SIZE option for clGetKernelWorkGroupInfo.
v2: should return the max global work size instead of current work size.
This funtion need return CL_INVALID_VALUE if the device is not a custom
device or kernel is not a built-in kernel.
we have 3 kind of built-in kernels for 1d/2d/3d memories, the max global
work size are decided by the dimension and memory type.
the piglit fail is caused by calling NON built-in kernels, so need send
patch to piglit later.
Signed-off-by: Luo Xionghu <xionghu.luo@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
-rw-r--r-- | src/cl_device_id.c | 34 | ||||
-rw-r--r-- | src/cl_device_id.h | 3 | ||||
-rw-r--r-- | src/cl_gt_device.h | 3 | ||||
-rw-r--r-- | src/cl_kernel.h | 2 | ||||
-rw-r--r-- | utests/CMakeLists.txt | 1 | ||||
-rw-r--r-- | utests/builtin_kernel_max_global_size.cpp | 30 |
6 files changed, 73 insertions, 0 deletions
diff --git a/src/cl_device_id.c b/src/cl_device_id.c index a0f0c99..ee3f2b7 100644 --- a/src/cl_device_id.c +++ b/src/cl_device_id.c @@ -515,6 +515,22 @@ cl_device_get_version(cl_device_id device, cl_int *ver) #include "cl_kernel.h" #include "cl_program.h" +static int +cl_check_builtin_kernel_dimension(cl_kernel kernel, cl_device_id device) +{ + const char * n = cl_kernel_get_name(kernel); + const char * builtin_kernels_2d = "__cl_copy_image_2d_to_2d;__cl_copy_image_2d_to_buffer;__cl_copy_buffer_to_image_2d;__cl_fill_image_2d;__cl_fill_image_2d_array;"; + const char * builtin_kernels_3d = "__cl_copy_image_3d_to_2d;__cl_copy_image_2d_to_3d;__cl_copy_image_3d_to_3d;__cl_copy_image_3d_to_buffer;__cl_copy_buffer_to_image_3d;__cl_fill_image_3d"; + if (!strstr(device->built_in_kernels, n)){ + return 0; + }else if(strstr(builtin_kernels_2d, n)){ + return 2; + }else if(strstr(builtin_kernels_3d, n)){ + return 3; + }else + return 1; + +} LOCAL size_t cl_get_kernel_max_wg_sz(cl_kernel kernel) @@ -543,6 +559,7 @@ cl_get_kernel_workgroup_info(cl_kernel kernel, size_t* param_value_size_ret) { int err = CL_SUCCESS; + int dimension = 0; if (UNLIKELY(device != &intel_ivb_gt1_device && device != &intel_ivb_gt2_device && device != &intel_baytrail_t_device && @@ -573,6 +590,23 @@ cl_get_kernel_workgroup_info(cl_kernel kernel, } DECL_FIELD(COMPILE_WORK_GROUP_SIZE, kernel->compile_wg_sz) DECL_FIELD(PRIVATE_MEM_SIZE, kernel->stack_size) + case CL_KERNEL_GLOBAL_WORK_SIZE: + dimension = cl_check_builtin_kernel_dimension(kernel, device); + if ( !dimension ) return CL_INVALID_VALUE; + if (param_value_size_ret != NULL) + *param_value_size_ret = sizeof(device->max_1d_global_work_sizes); + if (param_value) { + if (dimension == 1) { + memcpy(param_value, device->max_1d_global_work_sizes, sizeof(device->max_1d_global_work_sizes)); + }else if(dimension == 2){ + memcpy(param_value, device->max_2d_global_work_sizes, sizeof(device->max_2d_global_work_sizes)); + }else if(dimension == 3){ + memcpy(param_value, device->max_3d_global_work_sizes, sizeof(device->max_3d_global_work_sizes)); + }else + return CL_INVALID_VALUE; + + return CL_SUCCESS; + } default: return CL_INVALID_VALUE; }; diff --git a/src/cl_device_id.h b/src/cl_device_id.h index c4f8227..31bce47 100644 --- a/src/cl_device_id.h +++ b/src/cl_device_id.h @@ -30,6 +30,9 @@ struct _cl_device_id { cl_uint max_work_item_dimensions; // should be 3. size_t max_work_item_sizes[3]; // equal to maximum work group size. size_t max_work_group_size; // maximum work group size under simd16 mode. + size_t max_1d_global_work_sizes[3]; // maximum 1d global work size for builtin kernels. + size_t max_2d_global_work_sizes[3]; // maximum 2d global work size for builtin kernels. + size_t max_3d_global_work_sizes[3]; // maximum 3d global work size for builtin kernels. cl_uint preferred_vector_width_char; cl_uint preferred_vector_width_short; cl_uint preferred_vector_width_int; diff --git a/src/cl_gt_device.h b/src/cl_gt_device.h index 33ef1f0..3cd54eb 100644 --- a/src/cl_gt_device.h +++ b/src/cl_gt_device.h @@ -21,6 +21,9 @@ .device_type = CL_DEVICE_TYPE_GPU, .vendor_id = 0, /* == device_id (set when requested) */ .max_work_item_dimensions = 3, +.max_1d_global_work_sizes = {1024 * 1024 * 256, 1, 1}, +.max_2d_global_work_sizes = {8192, 8192, 1}, +.max_3d_global_work_sizes = {8192, 8192, 2048}, .preferred_vector_width_char = 8, .preferred_vector_width_short = 8, .preferred_vector_width_int = 4, diff --git a/src/cl_kernel.h b/src/cl_kernel.h index f4ed8d3..85a997d 100644 --- a/src/cl_kernel.h +++ b/src/cl_kernel.h @@ -59,6 +59,8 @@ struct _cl_kernel { cl_ulong local_mem_sz; /* local memory size specified in kernel args. */ size_t compile_wg_sz[3]; /* Required workgroup size by __attribute__((reqd_work_gro up_size(X, Y, Z))) qualifier.*/ + size_t global_work_sz[3]; /* maximum global size that can be used to execute a kernel + (i.e. global_work_size argument to clEnqueueNDRangeKernel.)*/ size_t stack_size; /* stack size per work item. */ cl_argument *args; /* To track argument setting */ uint32_t arg_n:31; /* Number of arguments */ diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index b30e6f9..034f112 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -181,6 +181,7 @@ set (utests_sources test_printf.cpp enqueue_fill_buf.cpp enqueue_built_in_kernels.cpp + builtin_kernel_max_global_size.cpp image_1D_buffer.cpp compare_image_2d_and_1d_array.cpp compiler_constant_expr.cpp diff --git a/utests/builtin_kernel_max_global_size.cpp b/utests/builtin_kernel_max_global_size.cpp new file mode 100644 index 0000000..c777564 --- /dev/null +++ b/utests/builtin_kernel_max_global_size.cpp @@ -0,0 +1,30 @@ +#include "utest_helper.hpp" + +void builtin_kernel_max_global_size(void) +{ + char* built_in_kernel_names; + size_t built_in_kernels_size; + cl_int err = CL_SUCCESS; + size_t ret_sz; + + + OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, 0, 0, &built_in_kernels_size); + built_in_kernel_names = (char* )malloc(built_in_kernels_size * sizeof(char) ); + OCL_CALL(clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, built_in_kernels_size, (void*)built_in_kernel_names, &ret_sz); + OCL_ASSERT(ret_sz == built_in_kernels_size); + cl_program built_in_prog = clCreateProgramWithBuiltInKernels(ctx, 1, &device, built_in_kernel_names, &err); + OCL_ASSERT(built_in_prog != NULL); + cl_kernel builtin_kernel_1d = clCreateKernel(built_in_prog, "__cl_copy_region_unalign_src_offset", &err); + OCL_ASSERT(builtin_kernel_1d != NULL); + size_t param_value_size; + void* param_value; + clGetKernelWorkGroupInfo(builtin_kernel_1d, device, CL_KERNEL_GLOBAL_WORK_SIZE, 0, NULL, ¶m_value_size); + param_value = malloc(param_value_size); + clGetKernelWorkGroupInfo(builtin_kernel_1d, device, CL_KERNEL_GLOBAL_WORK_SIZE, param_value_size, param_value, 0); + OCL_ASSERT(*(size_t*)param_value == 256 * 1024 *1024); + clReleaseKernel(builtin_kernel_1d); + clReleaseProgram(built_in_prog); + free(param_value); +} + +MAKE_UTEST_FROM_FUNCTION(builtin_kernel_max_global_size); |