[Beignet] [PATCH 3/3] Add the test cases for 1D Image Array

[He Junyan]

for this patch set, the compare_image_2d_and_1d_array output:

########## x y is (30, 15), color1 is (2 1 30 15), color2 is (2 1 30 30)
########## x y is (31, 0), color1 is (2 1 31 0), color2 is (2 1 31 0)
########## x y is (31, 1), color1 is (2 1 31 1), color2 is (2 1 31 2)
########## x y is (31, 2), color1 is (2 1 31 2), color2 is (2 1 31 4)
########## x y is (31, 3), color1 is (2 1 31 3), color2 is (2 1 31 6)
########## x y is (31, 4), color1 is (2 1 31 4), color2 is (2 1 31 8)
########## x y is (31, 5), color1 is (2 1 31 5), color2 is (2 1 31 10)
########## x y is (31, 6), color1 is (2 1 31 6), color2 is (2 1 31 12)
########## x y is (31, 7), color1 is (2 1 31 7), color2 is (2 1 31 14)
########## x y is (31, 8), color1 is (2 1 31 8), color2 is (2 1 31 16)
########## x y is (31, 9), color1 is (2 1 31 9), color2 is (2 1 31 18)
########## x y is (31, 10), color1 is (2 1 31 10), color2 is (2 1 31 20)
########## x y is (31, 11), color1 is (2 1 31 11), color2 is (2 1 31 22)
########## x y is (31, 12), color1 is (2 1 31 12), color2 is (2 1 31 24)
########## x y is (31, 13), color1 is (2 1 31 13), color2 is (2 1 31 26)
########## x y is (31, 14), color1 is (2 1 31 14), color2 is (2 1 31 28)
########## x y is (31, 15), color1 is (2 1 31 15), color2 is (2 1 31 30)

color1 is the result of image2d_t and color2 is the result of
image1d_array_t.
The h of the image1d_array_t seems always twice of the image2d_t.
I can not find the problem by now, any idea?




On Tue, 2014-06-17 at 12:07 +0800, junyan.he at inbox.com wrote:
> From: Junyan He <junyan.he at linux.intel.com>
> 
> Signed-off-by: Junyan He <junyan.he at linux.intel.com>
> ---
>  kernels/compare_image_2d_and_1d_array.cl | 12 +++++
>  kernels/test_get_image_info_array.cl     | 25 ++++++++++
>  utests/CMakeLists.txt                    |  2 +
>  utests/compare_image_2d_and_1d_array.cpp | 78 ++++++++++++++++++++++++++++++++
>  utests/compiler_get_image_info_array.cpp | 64 ++++++++++++++++++++++++++
>  5 files changed, 181 insertions(+)
>  create mode 100644 kernels/compare_image_2d_and_1d_array.cl
>  create mode 100644 kernels/test_get_image_info_array.cl
>  create mode 100644 utests/compare_image_2d_and_1d_array.cpp
>  create mode 100644 utests/compiler_get_image_info_array.cpp
> 
> 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 0000000..ff25834
> --- /dev/null
> +++ b/kernels/compare_image_2d_and_1d_array.cl
> @@ -0,0 +1,12 @@
> +__kernel void
> +compare_image_2d_and_1d_array(image2d_t a1, image1d_array_t a2, sampler_t sampler)
> +{
> +  int2 coord;
> +  int4 color1;
> +  int4 color2;
> +  coord.x =  get_global_id(0);
> +  coord.y = get_global_id(1);
> +  color1 = read_imagei(a1, sampler, coord);
> +  color2 = read_imagei(a2, sampler, coord);
> +  printf("########## x y is (%d, %d), 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 0000000..333da77
> --- /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 f0e62e2..641a73b 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 0000000..f989049
> --- /dev/null
> +++ b/utests/compare_image_2d_and_1d_array.cpp
> @@ -0,0 +1,78 @@
> +#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));
> +
> +#if 1
> +  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 = j & 0x3f;
> +      char b = i & 0x3f;
> +      image_data2[w * j + i] = image_data1[w * j + i] = a << 24 | b << 16 | 0x1 << 8 | 0x2;
> +    }
> +  }
> +#endif
> +
> +  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);
> +
> +#if 0
> +  OCL_MAP_BUFFER_GTT(0);
> +  for (int j = 0; j < h; ++j) {
> +    for (int i = 0; i < w; i++) {
> +      char a = (j)& 0x3f;
> +      char b = (i)& 0x3f;
> +      ((uint32_t*)buf_data[0])[j * w + i] = a<<24 | a<<16 | a<< 8 | b;
> +    }
> +  }
> +  OCL_UNMAP_BUFFER_GTT(0);
> +#endif
> +
> +  // 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_CLAMP, CL_FILTER_NEAREST);
> +
> +  // 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_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 0000000..970877d
> --- /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);