[Beignet] [PATCH 3/3] fix utest failures on VPG code.
xionghu.luo at intel.com
xionghu.luo at intel.com
Wed Apr 13 09:30:51 UTC 2016
From: Luo Xionghu <xionghu.luo at intel.com>
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 <xionghu.luo at intel.com>
---
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
More information about the Beignet
mailing list