utests: refine image 1d buffer test case.

We need to test large image 1d buffer read and write testing.

Signed-off-by: Zhigang Gong <zhigang.gong@intel.com>
Reviewed-by: "Yang, Rong R" <rong.r.yang@intel.com>

2 files changed, 32 insertions, 53 deletions

diff --git a/kernels/image_1D_buffer.cl b/kernels/image_1D_buffer.cl
index e8e0a863..2c1da694 100644
--- a/kernels/image_1D_buffer.cl
+++ b/kernels/image_1D_buffer.cl
@@ -1,13 +1,7 @@
-__kernel void image_1D_buffer(image1d_buffer_t image1, image1d_t image2, sampler_t sampler, __global int *results)
+__kernel void image_1D_buffer(image1d_buffer_t image1, image1d_buffer_t image2)
 {
    int x = get_global_id(0);
-   int offset = x;
 
-   int4 col = read_imagei(image1, x);
-   int4 test = (col != read_imagei(image2, sampler, x));
-
-   if (test.x || test.y || test.z || test.w)
-      results[offset] = 0;
-   else
-      results[offset] = 1;
+   uint4 color = read_imageui(image1, x);
+   write_imageui(image2, x, color);
 }
diff --git a/utests/image_1D_buffer.cpp b/utests/image_1D_buffer.cpp
index d8d761f0..e2cfcde5 100644
--- a/utests/image_1D_buffer.cpp
+++ b/utests/image_1D_buffer.cpp
@@ -3,78 +3,63 @@
 
 void image_1D_buffer(void)
 {
-  size_t buffer_sz = 1024;
-  char *buf_content = (char *)malloc(buffer_sz * sizeof(char));
+  size_t buffer_sz = 65536;
+  char *buf_content = (char *)malloc(buffer_sz * sizeof(int));
   int error;
   cl_image_desc image_desc;
   cl_image_format image_format;
-  cl_sampler sampler;
-  cl_mem image1, image2;
   cl_mem ret_mem = NULL;
 
   OCL_CREATE_KERNEL("image_1D_buffer");
 
   for (int32_t i = 0; i < (int32_t)buffer_sz; ++i)
-    buf_content[i] = (rand() & 127);
+    buf_content[i] = (rand() & 0xFFFFFFFF);
 
-  cl_mem buff = clCreateBuffer(ctx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
-                                      buffer_sz, buf_content, &error);
-  OCL_ASSERT(error == CL_SUCCESS);
+  OCL_CREATE_BUFFER(buf[0], CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, buffer_sz * sizeof(int), buf_content);
+  OCL_CREATE_BUFFER(buf[1], CL_MEM_READ_WRITE, buffer_sz * sizeof(int), NULL);
 
   memset(&image_desc, 0x0, sizeof(cl_image_desc));
   memset(&image_format, 0x0, sizeof(cl_image_format));
 
   image_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
-  image_desc.image_row_pitch = buffer_sz;
-  image_desc.image_width = buffer_sz / sizeof(uint32_t); //assume rgba32
-  image_desc.buffer = buff;
-
-  image_format.image_channel_order = CL_RGBA;
-  image_format.image_channel_data_type = CL_UNSIGNED_INT8;
-
-  image1 = clCreateImage(ctx, CL_MEM_READ_ONLY, &image_format,
-                        &image_desc, NULL, &error );
-  OCL_ASSERT(error == CL_SUCCESS);
-
-  error = clGetImageInfo(image1, CL_IMAGE_BUFFER, sizeof(ret_mem), &ret_mem, NULL);
-  OCL_ASSERT(error == CL_SUCCESS);
-  OCL_ASSERT(ret_mem == buff);
+  image_desc.image_row_pitch = buffer_sz * sizeof(int);
+  image_desc.image_width = buffer_sz; //assume r32
+  image_desc.buffer = buf[0];
 
+  image_format.image_channel_order = CL_R;
+  image_format.image_channel_data_type = CL_UNSIGNED_INT32;
 
-  memset(&image_desc, 0x0, sizeof(cl_image_desc));
-  image_desc.image_type = CL_MEM_OBJECT_IMAGE1D;
-  image_desc.image_width = buffer_sz / sizeof(uint32_t);
-  image2 = clCreateImage(ctx, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
-                         &image_format, &image_desc, buf_content, &error);
+  // Create the source image1d_buffer.
+  OCL_CREATE_IMAGE(buf[2], CL_MEM_READ_ONLY, &image_format, &image_desc, NULL);
+  error = clGetImageInfo(buf[2], CL_IMAGE_BUFFER, sizeof(ret_mem), &ret_mem, NULL);
   OCL_ASSERT(error == CL_SUCCESS);
+  OCL_ASSERT(ret_mem == buf[0]);
 
-  // Create sampler to use
-  sampler = clCreateSampler(ctx, false, CL_ADDRESS_NONE, CL_FILTER_NEAREST, &error );
+  // Create the destination image1d_buffer.
+  image_desc.buffer = buf[1];
+  OCL_CREATE_IMAGE(buf[3], CL_MEM_READ_ONLY, &image_format, &image_desc, NULL);
+  error = clGetImageInfo(buf[3], CL_IMAGE_BUFFER, sizeof(ret_mem), &ret_mem, NULL);
   OCL_ASSERT(error == CL_SUCCESS);
+  OCL_ASSERT(ret_mem == buf[1]);
 
-  cl_mem result_buf = buf[0] = clCreateBuffer(ctx, 0, buffer_sz, NULL, &error);
-  OCL_ASSERT(error == CL_SUCCESS);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[2]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[3]);
 
-  OCL_SET_ARG(0, sizeof(cl_mem), &image1);
-  OCL_SET_ARG(1, sizeof(cl_mem), &image2);
-  OCL_SET_ARG(2, sizeof(sampler), &sampler);
-  OCL_SET_ARG(3, sizeof(cl_mem), &result_buf);
-
-  globals[0] = buffer_sz/sizeof(int32_t);
+  globals[0] = buffer_sz;
   locals[0] = 16;
 
   OCL_NDRANGE(1);
 
   /* Now check the result. */
   OCL_MAP_BUFFER(0);
-  for (uint32_t i = 0; i < buffer_sz/sizeof(int32_t); i++)
-    OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 1);
+  OCL_MAP_BUFFER(1);
+  for (uint32_t i = 0; i < buffer_sz; i++) {
+    if (((uint32_t*)buf_data[1])[i] != ((uint32_t*)buf_data[0])[i])
+      printf("i %d expected %x got %x \n", i, ((uint32_t*)buf_data[0])[i], ((uint32_t*)buf_data[1])[i]);
+    OCL_ASSERT(((uint32_t*)buf_data[1])[i] == ((uint32_t*)buf_data[0])[i]);
+  }
   OCL_UNMAP_BUFFER(0);
-
-  clReleaseSampler(sampler);
-  clReleaseMemObject(image1);
-  clReleaseMemObject(image2);
-  clReleaseMemObject(buff);
+  OCL_UNMAP_BUFFER(1);
 }
 
 MAKE_UTEST_FROM_FUNCTION(image_1D_buffer);