fix failed cases for stand alone utest;

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@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

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 e7560101..e66359fc 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 43a804bc..ae088874 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 d678c0f0..4740c806 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 9601cab0..1fa9f0d7 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 457092f7..06807c20 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 094e019c..a2ec24a6 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 e6910cdb..d3e83735 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 1f076159..9f0adee3 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, 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 c2df7be7..8d706d05 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 a9dac2e1..491175db 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 bbff4353..832766e9 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 7c7dee3d..1bd2304f 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 9116608c..53a418fc 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 debdf940..1a4e074b 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 c960ae6b..c1496904 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 0c6992a8..1782df5c 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 793682b7..a5a2dda0 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 25f18c75..84029da7 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 426473af..77a19268 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 70b983b9..8ce77074 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)