Add the test cases for 1D Image Array

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

5 files changed, 183 insertions, 0 deletions

diff --git a/kernels/compare_image_2d_and_1d_array.cl b/kernels/compare_image_2d_and_1d_array.cl
new file mode 100644
index 00000000..6aabb43c
--- /dev/null
+++ b/kernels/compare_image_2d_and_1d_array.cl
@@ -0,0 +1,13 @@
+__kernel void
+compare_image_2d_and_1d_array(image2d_t a1, image1d_array_t a2, sampler_t sampler)
+{
+  float2 coord;
+  int4 color1;
+  int4 color2;
+  coord.x = (float)get_global_id(0) + 0.3f;
+  coord.y = (float)get_global_id(1) + 0.3f;
+  color1 = read_imagei(a1, sampler, coord);
+  color2 = read_imagei(a2, sampler, coord);
+//  printf("########## x y is (%f, %f), color1 is (%d %d %d %d), color2 is (%d %d %d %d)\n",
+//	  coord.x, coord.y, color1.x, color1.y, color1.z, color1.w, color2.x, color2.y, color2.z, color2.w);
+}
diff --git a/kernels/test_get_image_info_array.cl b/kernels/test_get_image_info_array.cl
new file mode 100644
index 00000000..333da776
--- /dev/null
+++ b/kernels/test_get_image_info_array.cl
@@ -0,0 +1,25 @@
+__kernel void
+test_get_image_info_array(__write_only image1d_array_t a1, __write_only image2d_array_t a2, __global int *result)
+{
+  int w, h, array_sz;
+
+  w = get_image_width(a1);
+  array_sz = (int)get_image_array_size(a1);
+  int channel_data_type = get_image_channel_data_type(a1);
+  int channel_order = get_image_channel_order(a1);
+  result[0] = w;
+  result[1] = array_sz;
+  result[2] = channel_data_type;
+  result[3] = channel_order;
+
+  w = get_image_width(a2);
+  h = get_image_height(a2);
+  array_sz = (int)get_image_array_size(a2);
+  channel_data_type = get_image_channel_data_type(a2);
+  channel_order = get_image_channel_order(a2);
+  result[4] = w;
+  result[5] = h;
+  result[6] = array_sz;
+  result[7] = channel_data_type;
+  result[8] = channel_order;
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index f0e62e21..641a73b1 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -122,6 +122,7 @@ set (utests_sources
   compiler_volatile.cpp
   compiler_copy_image1.cpp
   compiler_get_image_info.cpp
+  compiler_get_image_info_array.cpp
   compiler_vect_compare.cpp
   compiler_vector_load_store.cpp
   compiler_vector_inc.cpp
@@ -182,6 +183,7 @@ set (utests_sources
   enqueue_fill_buf.cpp
   enqueue_built_in_kernels.cpp
   image_1D_buffer.cpp
+  compare_image_2d_and_1d_array.cpp
   utest_assert.cpp
   utest.cpp
   utest_file_map.cpp
diff --git a/utests/compare_image_2d_and_1d_array.cpp b/utests/compare_image_2d_and_1d_array.cpp
new file mode 100644
index 00000000..f2c828ea
--- /dev/null
+++ b/utests/compare_image_2d_and_1d_array.cpp
@@ -0,0 +1,79 @@
+#include <string.h>
+#include "utest_helper.hpp"
+
+static void compare_image_2d_and_1d_array(void)
+{
+  const int w = 64;
+  const int h = 32;
+  cl_image_format format;
+  cl_image_desc desc;
+  cl_sampler sampler;
+
+  // Create the 1D array buffer.
+  memset(&desc, 0x0, sizeof(cl_image_desc));
+  memset(&format, 0x0, sizeof(cl_image_format));
+
+  uint32_t* image_data1 = (uint32_t *)malloc(w * h * sizeof(uint32_t));
+  uint32_t* image_data2 = (uint32_t *)malloc(w * h * sizeof(uint32_t));
+  for (int j = 0; j < h; j++) {
+    for (int i = 0; i < w; i++) {
+      char a = 0;
+      if (j % 2 == 0)
+        a = (j + 3) & 0x3f;
+
+      image_data2[w * j + i] = image_data1[w * j + i] = a << 24 | a << 16 | a << 8 | a;
+    }
+  }
+
+  format.image_channel_order = CL_RGBA;
+  format.image_channel_data_type = CL_UNSIGNED_INT8;
+  desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+  desc.image_width = w;
+  desc.image_height = h;
+  desc.image_row_pitch = w * sizeof(uint32_t);
+  OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data1);
+
+  // Create the 2D array buffer.
+  memset(&desc, 0x0, sizeof(cl_image_desc));
+  memset(&format, 0x0, sizeof(cl_image_format));
+
+  format.image_channel_order = CL_RGBA;
+  format.image_channel_data_type = CL_UNSIGNED_INT8;
+  desc.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
+  desc.image_width = w;
+  desc.image_array_size = h;
+  desc.image_row_pitch = w * sizeof(uint32_t);
+  OCL_CREATE_IMAGE(buf[1], CL_MEM_COPY_HOST_PTR, &format, &desc, image_data2);
+
+  OCL_CREATE_SAMPLER(sampler, CL_ADDRESS_REPEAT, CL_FILTER_LINEAR);
+
+  // Setup kernel and images
+  OCL_CREATE_KERNEL("compare_image_2d_and_1d_array");
+
+  // Run the kernel
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(cl_sampler), &sampler);
+  globals[0] = 32;
+  globals[1] = 16;
+  locals[0] = 32;
+  locals[1] = 16;
+  OCL_NDRANGE(2);
+
+  OCL_MAP_BUFFER_GTT(0);
+  OCL_MAP_BUFFER_GTT(1);
+  for (int j = 0; j < h; ++j) {
+    for (int i = 0; i < w; i++) {
+      // Because the array index will not join the sample caculation, the result should
+      // be different between the 2D and 1D_array.
+      if (j % 2 == 0)
+        OCL_ASSERT(((uint32_t*)buf_data[0])[j * w + i] == ((uint32_t*)buf_data[1])[j * w + i]);
+    }
+  }
+  OCL_UNMAP_BUFFER_GTT(0);
+  OCL_UNMAP_BUFFER_GTT(1);
+
+  OCL_CALL(clReleaseSampler, sampler);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compare_image_2d_and_1d_array);
diff --git a/utests/compiler_get_image_info_array.cpp b/utests/compiler_get_image_info_array.cpp
new file mode 100644
index 00000000..970877d1
--- /dev/null
+++ b/utests/compiler_get_image_info_array.cpp
@@ -0,0 +1,64 @@
+#include <string.h>
+#include "utest_helper.hpp"
+
+static void compiler_get_image_info_array(void)
+{
+  const int w = 256;
+  const int h = 512;
+  const int array_size1 = 10;
+  const int array_size2 = 3;
+  cl_image_format format;
+  cl_image_desc desc;
+
+  // Create the 1D array buffer.
+  memset(&desc, 0x0, sizeof(cl_image_desc));
+  memset(&format, 0x0, sizeof(cl_image_format));
+
+  format.image_channel_order = CL_RGBA;
+  format.image_channel_data_type = CL_UNSIGNED_INT8;
+  desc.image_type = CL_MEM_OBJECT_IMAGE1D_ARRAY;
+  desc.image_width = w;
+  desc.image_array_size = array_size1;
+  OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
+
+  // Create the 2D array buffer.
+  memset(&desc, 0x0, sizeof(cl_image_desc));
+  memset(&format, 0x0, sizeof(cl_image_format));
+
+  format.image_channel_order = CL_RGBA;
+  format.image_channel_data_type = CL_UNSIGNED_INT8;
+  desc.image_type = CL_MEM_OBJECT_IMAGE2D_ARRAY;
+  desc.image_width = w;
+  desc.image_height = h;
+  desc.image_array_size = array_size2;
+  OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL);
+
+  // Setup kernel and images
+  OCL_CREATE_KERNEL("test_get_image_info_array");
+
+  OCL_CREATE_BUFFER(buf[2], 0, 32 * sizeof(int), NULL);
+
+  // Run the kernel
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+  OCL_SET_ARG(2, sizeof(cl_mem), &buf[2]);
+  globals[0] = 32;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+
+  // Check result
+  OCL_MAP_BUFFER(2);
+  OCL_ASSERT(((int*)buf_data[2])[0] == w);
+  OCL_ASSERT(((int*)buf_data[2])[1] == array_size1);
+  OCL_ASSERT(((int*)buf_data[2])[2] == CL_UNSIGNED_INT8);
+  OCL_ASSERT(((int*)buf_data[2])[3] == CL_RGBA);
+
+  OCL_ASSERT(((int*)buf_data[2])[4] == w);
+  OCL_ASSERT(((int*)buf_data[2])[5] == h);
+  OCL_ASSERT(((int*)buf_data[2])[6] == array_size2);
+  OCL_ASSERT(((int*)buf_data[2])[7] == CL_UNSIGNED_INT8);
+  OCL_ASSERT(((int*)buf_data[2])[8] == CL_RGBA);
+  OCL_UNMAP_BUFFER(2);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_get_image_info_array);