[Beignet] [PATCH 4/5] fix failed cases for stand alone utest;
xionghu.luo at intel.com
xionghu.luo at intel.com
Thu Apr 21 10:50:40 UTC 2016
From: Luo Xionghu <xionghu.luo at intel.com>
1. use clEnqueueMapBuffer/Image instead of clEnqueueReadBuffer/Image;
2. add sanity check for clEnqueueMapImage;
v2: disable OpenCL 2.0 specific builtin cases for stand alone utest.
v3: don't hide failed cases. fix utest build warnings.
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/buildin_work_dim.cpp | 13 +++----------
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 | 8 +++++++-
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_clz.cpp | 16 ++++++++--------
utests/compiler_get_max_sub_group_size.cpp | 2 +-
utests/compiler_popcount.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 | 1 -
utests/utest_helper.hpp | 6 ++++--
20 files changed, 69 insertions(+), 95 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/buildin_work_dim.cpp b/utests/buildin_work_dim.cpp
index d678c0f..4740c80 100644
--- a/utests/buildin_work_dim.cpp
+++ b/utests/buildin_work_dim.cpp
@@ -3,8 +3,6 @@
static void buildin_work_dim(void)
{
// Setup kernel and buffers
-
- int result, err;
OCL_CREATE_KERNEL("buildin_work_dim");
OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int), NULL);
@@ -23,14 +21,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..1fa9f0d 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, 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..06807c2 100644
--- a/utests/builtin_global_linear_id.cpp
+++ b/utests/builtin_global_linear_id.cpp
@@ -31,7 +31,7 @@ static void builtin_global_linear_id(void)
{
// Setup kernel and buffers
- int dim, global_id[80], err, i, buf_len=1;
+ int dim, err, i, buf_len=1;
size_t offsets[3] = {0,0,0};
OCL_CREATE_KERNEL("builtin_global_linear_id");
@@ -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..d3e8373 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)
{
@@ -9,12 +10,17 @@ void builtin_kernel_max_global_size(void)
OCL_CALL (clGetDeviceInfo, device, CL_DEVICE_BUILT_IN_KERNELS, 0, 0, &built_in_kernels_size);
+ if(built_in_kernels_size == 0)
+ return;
+
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);
+ 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..8d706d0 100644
--- a/utests/builtin_local_linear_id.cpp
+++ b/utests/builtin_local_linear_id.cpp
@@ -32,7 +32,7 @@ static void builtin_local_linear_id(void)
{
// Setup kernel and buffers
- int dim, local_id[576], err, i, buf_len=1;
+ int dim, i, buf_len=1;
OCL_CREATE_KERNEL("builtin_local_linear_id");
OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_WRITE, sizeof(int)*576, NULL);
@@ -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_clz.cpp b/utests/compiler_clz.cpp
index 9116608..53a418f 100644
--- a/utests/compiler_clz.cpp
+++ b/utests/compiler_clz.cpp
@@ -81,13 +81,13 @@ void test(const char *kernel_name, int s_type)
{
for (uint32_t i = 0; i < n; ++i) {
if(sizeof(U) == 1 && i < 8 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
else if(sizeof(U) == 2 && i < 16 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
else if(sizeof(U) == 4 && i < 32 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
else if(sizeof(U) == 8 && i < 64 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i );
}
}
else // signed type
@@ -96,28 +96,28 @@ void test(const char *kernel_name, int s_type)
if(sizeof(U) == 1)
{
if( i < 8 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i+1 );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i+1 );
else if( i == 8 )
OCL_ASSERT(((U*)buf_data[1])[i] == 0 );
}
else if(sizeof(U) == 2)
{
if( i < 16 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i+1 );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i+1 );
else if( i == 16 )
OCL_ASSERT(((U*)buf_data[1])[i] == 0 );
}
else if(sizeof(U) == 4)
{
if( i < 32 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i+1 );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i+1 );
else if( i == 32 )
OCL_ASSERT(((U*)buf_data[1])[i] == 0 );
}
else if(sizeof(U) == 8)
{
if( i < 63 )
- OCL_ASSERT(((U*)buf_data[1])[i] == i+1 );
+ OCL_ASSERT(((U*)buf_data[1])[i] == (U)i+1 );
}
}
}
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_popcount.cpp b/utests/compiler_popcount.cpp
index c960ae6..c149690 100644
--- a/utests/compiler_popcount.cpp
+++ b/utests/compiler_popcount.cpp
@@ -51,7 +51,7 @@ void test(const char *kernel_name, int s_type)
OCL_MAP_BUFFER(1);
OCL_ASSERT(((T*)buf_data[1])[0] == 0);
for (int i = 1; i < n; ++i){
- OCL_ASSERT(((T*)buf_data[1])[i] == n-i-s_type);
+ OCL_ASSERT(((T*)buf_data[1])[i] == (T)n-i-s_type);
}
OCL_UNMAP_BUFFER(1);
}
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..a5a2dda 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] == (int)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..77a1926 100644
--- a/utests/utest_helper.cpp
+++ b/utests/utest_helper.cpp
@@ -209,7 +209,6 @@ clpanic(const char *msg, int rval)
char*
cl_do_kiss_path(const char *file, cl_device_id device)
{
- cl_int ver;
const char *sub_path = NULL;
char *ker_path = NULL;
const char *kiss_path = getenv("OCL_KERNEL_PATH");
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
More information about the Beignet
mailing list