[Beignet] [PATCH 3/3 V2] Add the test case for image 2d array fill

Zhigang Gong zhigang.gong at linux.intel.com
Wed Oct 22 18:51:51 PDT 2014


This version LGTM, just pushed, thanks.

On Wed, Oct 22, 2014 at 03:51:06PM +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/test_fill_image_2d_array.cl     | 13 +++++
>  utests/CMakeLists.txt                   |  1 +
>  utests/compiler_fill_image_2d_array.cpp | 84 +++++++++++++++++++++++++++++++++
>  3 files changed, 98 insertions(+)
>  create mode 100644 kernels/test_fill_image_2d_array.cl
>  create mode 100644 utests/compiler_fill_image_2d_array.cpp
> 
> diff --git a/kernels/test_fill_image_2d_array.cl b/kernels/test_fill_image_2d_array.cl
> new file mode 100644
> index 0000000..e756010
> --- /dev/null
> +++ b/kernels/test_fill_image_2d_array.cl
> @@ -0,0 +1,13 @@
> +__kernel void
> +test_fill_image_2d_array(__write_only image2d_array_t dst)
> +{
> +  int coordx;
> +  int coordy;
> +  int coordz;
> +  coordx = (int)get_global_id(0);
> +  coordy = (int)get_global_id(1);
> +  coordz = (int)get_global_id(2);
> +  uint4 color4 = {0, 1, 2 ,3};
> +  if (coordz < 7)
> +    write_imageui(dst, (int3)(coordx, coordy, coordz), color4);
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
> index 57d5909..cf3addc 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -188,6 +188,7 @@ set (utests_sources
>    image_1D_buffer.cpp
>    compare_image_2d_and_1d_array.cpp
>    compiler_fill_image_1d_array.cpp
> +  compiler_fill_image_2d_array.cpp
>    compiler_constant_expr.cpp
>    vload_bench.cpp
>    utest_assert.cpp
> diff --git a/utests/compiler_fill_image_2d_array.cpp b/utests/compiler_fill_image_2d_array.cpp
> new file mode 100644
> index 0000000..649b416
> --- /dev/null
> +++ b/utests/compiler_fill_image_2d_array.cpp
> @@ -0,0 +1,84 @@
> +#include <string.h>
> +#include "utest_helper.hpp"
> +
> +static void compiler_fill_image_2d_array(void)
> +{
> +  const size_t w = 64;
> +  const size_t h = 16;
> +  const size_t array = 8;
> +  cl_image_format format;
> +  cl_image_desc desc;
> +  size_t origin[3] = { };
> +  size_t region[3];
> +  uint32_t* dst;
> +
> +  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_row_pitch = 0;//w * sizeof(uint32_t);
> +  desc.image_array_size = array;
> +
> +  // Setup kernel and images
> +  OCL_CREATE_KERNEL("test_fill_image_2d_array");
> +
> +  OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
> +
> +  OCL_MAP_BUFFER_GTT(0);
> +  memset(buf_data[0], 0, sizeof(uint32_t) * w * h * array);
> +  OCL_UNMAP_BUFFER_GTT(0);
> +
> +  // Run the kernel
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> +  globals[0] = w/2;
> +  locals[0] = 16;
> +  globals[1] = h;
> +  locals[1] = 8;
> +  globals[2] = array;
> +  locals[2] = 8;
> +  OCL_NDRANGE(3);
> +
> +  // Check result
> +  region[0] = w;
> +  region[1] = h;
> +  region[2] = array;
> +  dst = (uint32_t*)malloc(w*h*array*sizeof(uint32_t));
> +  OCL_READ_IMAGE(buf[0], origin, region, dst);
> +
> +#if 0
> +  printf("------ The image result is: -------\n");
> +  for (uint32_t k = 0; k < array; k++) {
> +    for (uint32_t j = 0; j < h; j++) {
> +      for (uint32_t i = 0; i < w; i++) {
> +        printf(" %2x", dst[k*h*w + j*w + i]);
> +      }
> +      printf("\n");
> +    }
> +    printf("\n");
> +  }
> +#endif
> +
> +  for (uint32_t k = 0; k < array - 1; k++) {
> +    for (uint32_t j = 0; j < h; j++) {
> +      for (uint32_t i = 0; i < w/2; i++) {
> +        OCL_ASSERT(dst[k*w*h + j*w + i] == 0x03020100);
> +      }
> +      for (uint32_t i = w/2; i < w; i++) {
> +        OCL_ASSERT(dst[k*w*h + j*w + i] == 0);
> +      }
> +    }
> +  }
> +
> +  for (uint32_t j = 0; j < h; j++) {
> +    for (uint32_t i = 0; i < w; i++) {
> +      OCL_ASSERT(dst[(array - 1)*w*h + j*w + i] == 0x0);
> +    }
> +  }
> +  free(dst);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_2d_array);
> -- 
> 1.9.1
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list