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 -cl-std=CL2.0 for OpenCL 2.0 specific cases; 3. add sanity check for clEnqueueMapImage; 4. set NON strict conformance as default.
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 | 45 +++++++++++++++++------------- utests/buildin_work_dim.cpp | 11 ++------ utests/builtin_global_id.cpp | 16 ++++------- utests/builtin_global_linear_id.cpp | 16 ++++------- 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 | 16 ++++------- 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/compiler_workgroup_broadcast.cpp | 2 +- utests/compiler_workgroup_reduce.cpp | 14 +++++----- utests/runtime_alloc_host_ptr_buffer.cpp | 6 ++-- utests/utest_generator.py | 8 ++++-- utests/utest_helper.cpp | 2 +- utests/utest_helper.hpp | 16 +++++++++-- 21 files changed, 101 insertions(+), 112 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..71f29f0 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 @@ -178,7 +176,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 +186,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 +197,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 +219,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 +229,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 +238,37 @@ 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_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..49441b5 100644 --- a/utests/builtin_global_linear_id.cpp +++ b/utests/builtin_global_linear_id.cpp @@ -33,7 +33,7 @@ static void builtin_global_linear_id(void) // Setup kernel and buffers int dim, global_id[80], err, i, buf_len=1; size_t offsets[3] = {0,0,0}; - OCL_CREATE_KERNEL("builtin_global_linear_id"); + OCL_CREATE_KERNEL_20("builtin_global_linear_id"); OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int)*80, NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -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..e8af05f 100644 --- a/utests/builtin_local_linear_id.cpp +++ b/utests/builtin_local_linear_id.cpp @@ -33,7 +33,7 @@ static void builtin_local_linear_id(void) // Setup kernel and buffers int dim, local_id[576], err, i, buf_len=1; - OCL_CREATE_KERNEL("builtin_local_linear_id"); + OCL_CREATE_KERNEL_20("builtin_local_linear_id"); OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int)*576, NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -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/compiler_workgroup_broadcast.cpp b/utests/compiler_workgroup_broadcast.cpp index d45e5d8..debed09f 100644 --- a/utests/compiler_workgroup_broadcast.cpp +++ b/utests/compiler_workgroup_broadcast.cpp @@ -11,7 +11,7 @@ void compiler_workgroup_broadcast(void) uint32_t src[n]; // Setup kernel and buffers - OCL_CREATE_KERNEL("compiler_workgroup_broadcast"); + OCL_CREATE_KERNEL_20("compiler_workgroup_broadcast"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); diff --git a/utests/compiler_workgroup_reduce.cpp b/utests/compiler_workgroup_reduce.cpp index 4097843..0d984d1 100644 --- a/utests/compiler_workgroup_reduce.cpp +++ b/utests/compiler_workgroup_reduce.cpp @@ -9,7 +9,7 @@ void compiler_workgroup_reduce_min_uniform(void) uint32_t src = 253; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_min_uniform"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_min_uniform"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); OCL_SET_ARG(0, sizeof(uint32_t), &src); OCL_SET_ARG(1, sizeof(cl_mem), &buf[0]); @@ -41,7 +41,7 @@ void compiler_workgroup_reduce_min_uint(void) uint32_t* src = test_array_uint; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_min_uint"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_min_uint"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -72,7 +72,7 @@ void compiler_workgroup_reduce_max_uint(void) uint32_t* src = test_array_uint; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_max_uint"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_max_uint"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -103,7 +103,7 @@ void compiler_workgroup_reduce_add_uint(void) uint32_t* src = test_array_uint; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_add_uint"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_add_uint"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -151,7 +151,7 @@ void compiler_workgroup_reduce_min_float(void) float* src = test_array_float; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_min_float"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_min_float"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -182,7 +182,7 @@ void compiler_workgroup_reduce_max_float(void) float* src = test_array_float; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_max_float"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_max_float"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); @@ -213,7 +213,7 @@ void compiler_workgroup_reduce_add_float(void) float* src = test_array_float; // Setup kernel and buffers - OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_reduce", "compiler_workgroup_reduce_add_float"); + OCL_CREATE_KERNEL_FROM_FILE_20("compiler_workgroup_reduce", "compiler_workgroup_reduce_add_float"); OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL); OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL); OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); 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..7128168 100644 --- a/utests/utest_helper.hpp +++ b/utests/utest_helper.hpp @@ -66,6 +66,11 @@ extern EGLSurface eglSurface; OCL_CALL (cl_kernel_init, NAME".cl", NAME, SOURCE, NULL); \ } while (0) +#define OCL_CREATE_KERNEL_20(NAME) \ + do { \ + OCL_CALL (cl_kernel_init, NAME".cl", NAME, SOURCE, "-cl-std=CL2.0"); \ + } while (0) + #define OCL_DESTROY_KERNEL_KEEP_PROGRAM(KEEP_PROGRAM) \ do { \ cl_kernel_destroy(!(KEEP_PROGRAM)); \ @@ -76,6 +81,11 @@ extern EGLSurface eglSurface; OCL_CALL(cl_kernel_init, FILE_NAME".cl", KERNEL_NAME, SOURCE, NULL); \ } while (0) +#define OCL_CREATE_KERNEL_FROM_FILE_20(FILE_NAME, KERNEL_NAME) \ + do { \ + OCL_CALL(cl_kernel_init, FILE_NAME".cl", KERNEL_NAME, SOURCE, "-cl-std=CL2.0"); \ + } while (0) + #define OCL_FLUSH() \ do { \ OCL_CALL(clFlush, queue); \ @@ -129,7 +139,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 +172,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
