From: Luo Xionghu <[email protected]> use NOT_BUILD_STAND_ALONE_UTEST to disable the cases failed on VPG now. 1. use clEnqueueMapBuffer/Image instead of clEnqueueReadBuffer/Image; 2. add sanity check for clEnqueueMapImage; 3. set NON strict conformance as default.
v2: disable OpenCL 2.0 specific builtin cases for stand alone utest. Signed-off-by: Luo Xionghu <[email protected]> --- kernels/test_fill_image_2d_array.cl | 2 +- kernels/test_get_arg_info.cl | 2 +- utests/CMakeLists.txt | 53 ++++++++++++++++-------------- utests/buildin_work_dim.cpp | 11 ++----- utests/builtin_global_id.cpp | 16 +++------ utests/builtin_global_linear_id.cpp | 14 +++----- utests/builtin_global_size.cpp | 9 ++--- utests/builtin_kernel_max_global_size.cpp | 5 ++- utests/builtin_local_id.cpp | 16 +++------ utests/builtin_local_linear_id.cpp | 14 +++----- utests/builtin_local_size.cpp | 10 ++---- utests/builtin_num_groups.cpp | 10 ++---- utests/compiler_cl_finish.cpp | 1 + utests/compiler_get_max_sub_group_size.cpp | 2 +- utests/compiler_unstructured_branch3.cpp | 4 +++ utests/runtime_alloc_host_ptr_buffer.cpp | 6 ++-- utests/utest_generator.py | 8 +++-- utests/utest_helper.cpp | 2 +- utests/utest_helper.hpp | 6 ++-- 19 files changed, 85 insertions(+), 106 deletions(-) diff --git a/kernels/test_fill_image_2d_array.cl b/kernels/test_fill_image_2d_array.cl index e756010..e66359f 100644 --- a/kernels/test_fill_image_2d_array.cl +++ b/kernels/test_fill_image_2d_array.cl @@ -9,5 +9,5 @@ test_fill_image_2d_array(__write_only image2d_array_t dst) coordz = (int)get_global_id(2); uint4 color4 = {0, 1, 2 ,3}; if (coordz < 7) - write_imageui(dst, (int3)(coordx, coordy, coordz), color4); + write_imageui(dst, (int4)(coordx, coordy, coordz, 0), color4); } diff --git a/kernels/test_get_arg_info.cl b/kernels/test_get_arg_info.cl index 43a804b..ae08887 100644 --- a/kernels/test_get_arg_info.cl +++ b/kernels/test_get_arg_info.cl @@ -3,6 +3,6 @@ typedef struct _test_arg_struct { int b; }test_arg_struct; -kernel void test_get_arg_info(read_only global float const volatile *src, read_write local int read_only *dst, test_arg_struct extra) { +kernel void test_get_arg_info(read_only global float const volatile *src, read_write local int *dst, test_arg_struct extra) { } diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index f4e85fb..d57b0ee 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -132,7 +132,6 @@ set (utests_sources compiler_rhadd.cpp compiler_rotate.cpp compiler_saturate.cpp - compiler_saturate_sub.cpp compiler_shift_right.cpp compiler_short_scatter.cpp compiler_smoothstep.cpp @@ -143,7 +142,6 @@ set (utests_sources compiler_uint3_unaligned_copy.cpp compiler_upsample_int.cpp compiler_upsample_long.cpp - compiler_unstructured_branch0.cpp compiler_unstructured_branch1.cpp compiler_unstructured_branch2.cpp compiler_unstructured_branch3.cpp @@ -156,8 +154,6 @@ set (utests_sources compiler_math.cpp compiler_atomic_functions.cpp compiler_async_copy.cpp - compiler_workgroup_broadcast.cpp - compiler_workgroup_reduce.cpp compiler_async_stride_copy.cpp compiler_insn_selection_min.cpp compiler_insn_selection_max.cpp @@ -178,7 +174,6 @@ set (utests_sources compiler_vector_load_store.cpp compiler_vector_inc.cpp compiler_cl_finish.cpp - get_cl_info.cpp builtin_atan2.cpp builtin_bitselect.cpp builtin_frexp.cpp @@ -189,9 +184,6 @@ set (utests_sources builtin_shuffle.cpp builtin_shuffle2.cpp builtin_sign.cpp - builtin_lgamma.cpp - builtin_lgamma_r.cpp - builtin_tgamma.cpp buildin_work_dim.cpp builtin_global_size.cpp builtin_local_size.cpp @@ -203,12 +195,8 @@ set (utests_sources builtin_exp.cpp builtin_convert_sat.cpp sub_buffer.cpp - runtime_createcontext.cpp runtime_set_kernel_arg.cpp runtime_null_kernel_arg.cpp - runtime_event.cpp - runtime_barrier_list.cpp - runtime_marker_list.cpp runtime_compile_link.cpp compiler_long.cpp compiler_long_2.cpp @@ -229,9 +217,6 @@ set (utests_sources compiler_private_const.cpp compiler_private_data_overflow.cpp compiler_getelementptr_bitcast.cpp - compiler_sub_group_any.cpp - compiler_sub_group_all.cpp - compiler_time_stamp.cpp compiler_double_precision.cpp compiler_double.cpp compiler_double_div.cpp @@ -242,14 +227,8 @@ set (utests_sources profiling_exec.cpp enqueue_copy_buf.cpp enqueue_copy_buf_unaligned.cpp - test_printf.cpp enqueue_fill_buf.cpp - builtin_kernel_max_global_size.cpp - image_1D_buffer.cpp image_from_buffer.cpp - compare_image_2d_and_1d_array.cpp - compiler_fill_image_1d_array.cpp - compiler_fill_image_2d_array.cpp compiler_constant_expr.cpp compiler_assignment_operation_in_if.cpp vload_bench.cpp @@ -257,13 +236,39 @@ set (utests_sources runtime_alloc_host_ptr_buffer.cpp runtime_use_host_ptr_image.cpp compiler_get_max_sub_group_size.cpp - compiler_get_sub_group_local_id.cpp compiler_sub_group_shuffle.cpp - builtin_global_linear_id.cpp - builtin_local_linear_id.cpp compiler_mix.cpp compiler_bsort.cpp) +if (NOT_BUILD_STAND_ALONE_UTEST) + SET(utests_sources + ${utests_sources} + compiler_saturate_sub.cpp + compiler_unstructured_branch0.cpp + get_cl_info.cpp + builtin_lgamma.cpp + builtin_lgamma_r.cpp + builtin_tgamma.cpp + runtime_createcontext.cpp + runtime_event.cpp + runtime_barrier_list.cpp + runtime_marker_list.cpp + compiler_sub_group_any.cpp + compiler_sub_group_all.cpp + compiler_time_stamp.cpp + test_printf.cpp + builtin_kernel_max_global_size.cpp + image_1D_buffer.cpp + compare_image_2d_and_1d_array.cpp + compiler_fill_image_1d_array.cpp + compiler_fill_image_2d_array.cpp + compiler_workgroup_broadcast.cpp + compiler_workgroup_reduce.cpp + builtin_global_linear_id.cpp + builtin_local_linear_id.cpp + compiler_get_sub_group_local_id.cpp) +endif (NOT_BUILD_STAND_ALONE_UTEST) + if (LLVM_VERSION_NODOT VERSION_GREATER 34) SET(utests_sources ${utests_sources} diff --git a/utests/buildin_work_dim.cpp b/utests/buildin_work_dim.cpp index d678c0f..f48b946 100644 --- a/utests/buildin_work_dim.cpp +++ b/utests/buildin_work_dim.cpp @@ -23,14 +23,9 @@ static void buildin_work_dim(void) // Run the kernel OCL_NDRANGE(i); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int), &result, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - - OCL_ASSERT( result == i); + OCL_MAP_BUFFER(0); + OCL_ASSERT( ((int*)buf_data[0])[0]== i); + OCL_UNMAP_BUFFER(0); } } diff --git a/utests/builtin_global_id.cpp b/utests/builtin_global_id.cpp index 9601cab..6a1f644 100644 --- a/utests/builtin_global_id.cpp +++ b/utests/builtin_global_id.cpp @@ -28,7 +28,7 @@ static void builtin_global_id(void) { // Setup kernel and buffers - int dim, global_id[80], err, i, buf_len=1; + int dim, err, i, buf_len=1; OCL_CREATE_KERNEL("builtin_global_id"); OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int)*80, NULL); @@ -53,24 +53,18 @@ static void builtin_global_id(void) OCL_NDRANGE( dim ); clFinish(queue); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * buf_len, &global_id, 0, NULL, NULL); - - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - + OCL_MAP_BUFFER(0); #if udebug for(i = 0; i < buf_len; i++) { - printf("%2d ", global_id[i]); + printf("%2d ", ((int*)buf_data[0])[i]); if ((i + 1) % 3 == 0) printf("\n"); } #endif for( i = 0; i < buf_len; i++) - OCL_ASSERT( global_id[i] == i); + OCL_ASSERT( ((int*)buf_data[0])[i] == i); + OCL_UNMAP_BUFFER(0); } } diff --git a/utests/builtin_global_linear_id.cpp b/utests/builtin_global_linear_id.cpp index 457092f..d157b63 100644 --- a/utests/builtin_global_linear_id.cpp +++ b/utests/builtin_global_linear_id.cpp @@ -65,24 +65,18 @@ static void builtin_global_linear_id(void) clFinish(queue); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * buf_len, &global_id, 0, NULL, NULL); - - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - + OCL_MAP_BUFFER(0); #if udebug for(i = 0; i < buf_len; i++) { - printf("%2d ", global_id[i]); + printf("%2d ", ((int*)buf_data[0])[i]); if ((i + 1) % 3 == 0) printf("\n"); } #endif for( i = 0; i < buf_len; i++) - OCL_ASSERT( global_id[i] == i); + OCL_ASSERT( ((int*)buf_data[0])[i] == i); + OCL_UNMAP_BUFFER(0); } } diff --git a/utests/builtin_global_size.cpp b/utests/builtin_global_size.cpp index 094e019..a2ec24a 100644 --- a/utests/builtin_global_size.cpp +++ b/utests/builtin_global_size.cpp @@ -80,12 +80,8 @@ static void builtin_global_size(void) // Run the kernel OCL_NDRANGE( dim ); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int), &global_size, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } + OCL_MAP_BUFFER(0); + global_size = ((int*)buf_data[0])[0]; //printf("get_global_size(%d) = %d (dimension:%d)\n", dim_arg_global, global_size, dim); @@ -101,6 +97,7 @@ static void builtin_global_size(void) OCL_ASSERT( global_size == 1); #endif } + OCL_UNMAP_BUFFER(0); } } } diff --git a/utests/builtin_kernel_max_global_size.cpp b/utests/builtin_kernel_max_global_size.cpp index e6910cd..4bd57f6 100644 --- a/utests/builtin_kernel_max_global_size.cpp +++ b/utests/builtin_kernel_max_global_size.cpp @@ -1,4 +1,5 @@ #include "utest_helper.hpp" +#include <string.h> void builtin_kernel_max_global_size(void) { @@ -14,7 +15,9 @@ void builtin_kernel_max_global_size(void) 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); + char* first_kernel = strtok(built_in_kernel_names, ";"); + OCL_ASSERT(first_kernel); + cl_kernel builtin_kernel_1d = clCreateKernel(built_in_prog, first_kernel, &err); OCL_ASSERT(builtin_kernel_1d != NULL); size_t param_value_size; void* param_value; diff --git a/utests/builtin_local_id.cpp b/utests/builtin_local_id.cpp index 1f07615..3a93f91 100644 --- a/utests/builtin_local_id.cpp +++ b/utests/builtin_local_id.cpp @@ -32,7 +32,7 @@ static void builtin_local_id(void) { // Setup kernel and buffers - int dim, local_id[576], err, i, buf_len=1; + int dim, err, i, buf_len=1; OCL_CREATE_KERNEL("builtin_local_id"); OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int)*576, NULL); @@ -57,24 +57,18 @@ static void builtin_local_id(void) OCL_NDRANGE( dim ); clFinish(queue); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * buf_len, &local_id, 0, NULL, NULL); - - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - + OCL_MAP_BUFFER(0); #if udebug for(i = 0; i < buf_len; i++) { - printf("%2d ", local_id[i]); + printf("%2d ", ((int*)buf_data[0])[i]); if ((i + 1) % 4 == 0) printf("\n"); } #endif for( i = 0; i < buf_len; i++) - OCL_ASSERT( local_id[i] == i); + OCL_ASSERT( ((int*)buf_data[0])[i] == i); + OCL_UNMAP_BUFFER(0); } } diff --git a/utests/builtin_local_linear_id.cpp b/utests/builtin_local_linear_id.cpp index c2df7be..e485708 100644 --- a/utests/builtin_local_linear_id.cpp +++ b/utests/builtin_local_linear_id.cpp @@ -57,24 +57,18 @@ static void builtin_local_linear_id(void) OCL_NDRANGE( dim ); clFinish(queue); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int) * buf_len, &local_id, 0, NULL, NULL); - - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - + OCL_MAP_BUFFER(0); #if udebug for(i = 0; i < buf_len; i++) { - printf("%2d ", local_id[i]); + printf("%2d ", ((int*)buf_data[0])[i]); if ((i + 1) % 4 == 0) printf("\n"); } #endif for( i = 0; i < buf_len; i++) - OCL_ASSERT( local_id[i] == i); + OCL_ASSERT( ((int*)buf_data[0])[i] == i); + OCL_UNMAP_BUFFER(0); } } diff --git a/utests/builtin_local_size.cpp b/utests/builtin_local_size.cpp index a9dac2e..491175d 100644 --- a/utests/builtin_local_size.cpp +++ b/utests/builtin_local_size.cpp @@ -65,13 +65,8 @@ static void builtin_local_size(void) // Run the kernel OCL_NDRANGE( dim ); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int), &local_size, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - + OCL_MAP_BUFFER(0); + local_size = ((int*)buf_data[0])[0]; #if udebug printf("get_local_size(%d) = %d (dimension:%d)\n", dim_arg_global, local_size, dim); #endif @@ -81,6 +76,7 @@ static void builtin_local_size(void) { OCL_ASSERT( local_size == 1); } + OCL_UNMAP_BUFFER(0); } } } diff --git a/utests/builtin_num_groups.cpp b/utests/builtin_num_groups.cpp index bbff435..832766e 100644 --- a/utests/builtin_num_groups.cpp +++ b/utests/builtin_num_groups.cpp @@ -62,13 +62,8 @@ static void builtin_num_groups(void) // Run the kernel OCL_NDRANGE( dim ); - err = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(int), &num_groups, 0, NULL, NULL); - if (err != CL_SUCCESS) - { - printf("Error: Failed to read output array! %d\n", err); - exit(1); - } - + OCL_MAP_BUFFER(0); + num_groups = ((int*)buf_data[0])[0]; #if udebug printf("get_num_groups(%d) = %d (dimension:%d)\n", dim_arg_global, num_groups, dim); #endif @@ -78,6 +73,7 @@ static void builtin_num_groups(void) { OCL_ASSERT( num_groups == 1); } + OCL_UNMAP_BUFFER(0); } } } diff --git a/utests/compiler_cl_finish.cpp b/utests/compiler_cl_finish.cpp index 7c7dee3..1bd2304 100644 --- a/utests/compiler_cl_finish.cpp +++ b/utests/compiler_cl_finish.cpp @@ -34,6 +34,7 @@ static void compiler_cl_finish(void) T_GET(t1); OCL_MAP_BUFFER(0); T_GET(t2); + OCL_UNMAP_BUFFER(0); t_map_w_fin = T_LAPSE(t1, t2); // 2nd time map without clFinish diff --git a/utests/compiler_get_max_sub_group_size.cpp b/utests/compiler_get_max_sub_group_size.cpp index debdf94..1a4e074 100644 --- a/utests/compiler_get_max_sub_group_size.cpp +++ b/utests/compiler_get_max_sub_group_size.cpp @@ -24,7 +24,7 @@ void compiler_get_max_sub_group_size(void) OCL_MAP_BUFFER(0); int* dst = (int *)buf_data[0]; for (int32_t i = 0; i < (int32_t) n; ++i){ - OCL_ASSERT(8 == dst[i] || 16 == dst[i]); + OCL_ASSERT(8 == dst[i] || 16 == dst[i] || 32 == dst[i]); } OCL_UNMAP_BUFFER(0); } diff --git a/utests/compiler_unstructured_branch3.cpp b/utests/compiler_unstructured_branch3.cpp index 0c6992a..1782df5 100644 --- a/utests/compiler_unstructured_branch3.cpp +++ b/utests/compiler_unstructured_branch3.cpp @@ -37,6 +37,8 @@ static void compiler_unstructured_branch3(void) OCL_MAP_BUFFER(1); for (uint32_t i = 0; i < n; ++i) OCL_ASSERT(((uint32_t*)buf_data[1])[i] == 3); + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); // Third control flow OCL_MAP_BUFFER(0); @@ -52,6 +54,8 @@ static void compiler_unstructured_branch3(void) OCL_ASSERT(((int32_t*)buf_data[1])[i] == 2); for (uint32_t i = 8; i < n; ++i) OCL_ASSERT(((int32_t*)buf_data[1])[i] == 3); + OCL_UNMAP_BUFFER(0); + OCL_UNMAP_BUFFER(1); } MAKE_UTEST_FROM_FUNCTION(compiler_unstructured_branch3); diff --git a/utests/runtime_alloc_host_ptr_buffer.cpp b/utests/runtime_alloc_host_ptr_buffer.cpp index 793682b..a1866a7 100644 --- a/utests/runtime_alloc_host_ptr_buffer.cpp +++ b/utests/runtime_alloc_host_ptr_buffer.cpp @@ -16,10 +16,10 @@ static void runtime_alloc_host_ptr_buffer(void) OCL_NDRANGE(1); // Check result - uint32_t* mapptr = (uint32_t*)clEnqueueMapBuffer(queue, buf[0], CL_TRUE, CL_MAP_READ, 0, n*sizeof(uint32_t), 0, NULL, NULL, NULL); + OCL_MAP_BUFFER(0); for (uint32_t i = 0; i < n; ++i) - OCL_ASSERT(mapptr[i] == i / 2); - clEnqueueUnmapMemObject(queue, buf[0], mapptr, 0, NULL, NULL); + OCL_ASSERT(((int*)buf_data[0])[i] == i / 2); + OCL_UNMAP_BUFFER(0); } MAKE_UTEST_FROM_FUNCTION(runtime_alloc_host_ptr_buffer); diff --git a/utests/utest_generator.py b/utests/utest_generator.py index 91cc938..9110c99 100644 --- a/utests/utest_generator.py +++ b/utests/utest_generator.py @@ -361,11 +361,15 @@ static void %s_%s(void) funcrun=''' // Run the kernel: + //int errRead = clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(%s) * count_input, gpu_data, 0, NULL, NULL); OCL_NDRANGE( 1 ); - clEnqueueReadBuffer( queue, buf[0], CL_TRUE, 0, sizeof(%s) * count_input, gpu_data, 0, NULL, NULL); -'''%(self.inputtype.__len__()+1) + OCL_MAP_BUFFER(0); +'''%(self.argtype(0,index)) funcline += [ funcrun ] + text = ''' memcpy(gpu_data, buf_data[0], sizeof(gpu_data)); ''' + funcline += [ text ] + funcsprintfa=' sprintf(log, \"' funcsprintfb='' if (self.returnVector(index) == 1 and self.argvector(0,index) != 1): diff --git a/utests/utest_helper.cpp b/utests/utest_helper.cpp index 426473a..db0ac51 100644 --- a/utests/utest_helper.cpp +++ b/utests/utest_helper.cpp @@ -808,7 +808,7 @@ float select_ulpsize(float ULPSIZE_FAST_MATH, float ULPSIZE_NO_FAST_MATH) const char* env_strict = getenv("OCL_STRICT_CONFORMANCE"); float ULPSIZE_FACTOR = ULPSIZE_NO_FAST_MATH; - if (env_strict != NULL && strcmp(env_strict, "0") == 0 ) + if (env_strict == NULL || strcmp(env_strict, "0") == 0 ) ULPSIZE_FACTOR = ULPSIZE_FAST_MATH; return ULPSIZE_FACTOR; diff --git a/utests/utest_helper.hpp b/utests/utest_helper.hpp index 70b983b..8ce7707 100644 --- a/utests/utest_helper.hpp +++ b/utests/utest_helper.hpp @@ -129,7 +129,7 @@ extern EGLSurface eglSurface; size_t size = 0; \ status = clGetMemObjectInfo(buf[ID], CL_MEM_SIZE, sizeof(size), &size, NULL);\ if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ - RET = FN(__VA_ARGS__, CL_TRUE, CL_MAP_READ, 0, size, 0, NULL, NULL, &status);\ + RET = FN(__VA_ARGS__, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, 0, size, 0, NULL, NULL, &status);\ if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ } while (0) @@ -162,9 +162,11 @@ extern EGLSurface eglSurface; size_t image_depth= 0; \ status = clGetImageInfo(buf[ID], CL_IMAGE_DEPTH, sizeof(image_depth), &image_depth, NULL);\ if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ + if(image_depth == 0) image_depth = 1; \ + if(image_height == 0) image_height = 1; \ size_t origin[3] = {0, 0, 0}; \ size_t region[3] = {image_width, image_height, image_depth}; \ - RET = FN(__VA_ARGS__, CL_TRUE, CL_MAP_READ, origin, region, &image_row_pitch, &image_slice_pitch, 0, NULL, NULL, &status);\ + RET = FN(__VA_ARGS__, CL_TRUE, CL_MAP_READ|CL_MAP_WRITE, origin, region, &image_row_pitch, &image_slice_pitch, 0, NULL, NULL, &status);\ if (status != CL_SUCCESS) OCL_THROW_ERROR(FN, status); \ } while (0) -- 2.1.4 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
