[Beignet] [PATCH 5/5] test cases for image3d_t

Xing, Homer homer.xing at intel.com
Tue May 7 17:25:21 PDT 2013


Try again?

These patches works very well on my machine. All test cases passed.

Homer

-----Original Message-----
From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com] 
Sent: Tuesday, May 07, 2013 3:13 PM
To: Xing, Homer
Cc: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH 5/5] test cases for image3d_t

Homer,

I met the following errors when I tried to run the utest after applied this patchset.

compiler_copy_image:
utest_run: /home/gong/git/fdo/beignet/backend/src/backend/gen_encoder.cpp:231: void gbe::GenEncoder::setSrc0(GenInstruction*, gbe::GenRegister): Assertion `reg.nr < 128' failed.
Aborted (core dumped)

compiler_copy_image_3d:
  compiler_copy_image_3d()    [FAILED]
    Error: ((uint32_t*)buf_data[0])[k*w*h + j*w + i] == ((uint32_t*)buf_data[1])[k*w*h + j*w + i]
  at file /home/gong/git/fdo/beignet/utests/compiler_copy_image_3d.cpp, function compiler_copy_image_3d, line 50

compiler_fill_image:
  compiler_fill_image()    [FAILED]
    Error: ((uint32_t*)buf_data[0])[j * w + i] == 0x78563412
  at file /home/gong/git/fdo/beignet/utests/compiler_fill_image.cpp, function compiler_fill_image, line 36

compiler_fill_image_3d:
  compiler_fill_image_3d()    [FAILED]
    Error: ((uint32_t*)buf_data[0])[k*w*h + j*w + i] == 0x78563412
  at file /home/gong/git/fdo/beignet/utests/compiler_fill_image_3d.cpp, function compiler_fill_image_3d, line 40

Couldy you take a look at this patchset? I just doubt maybe this is not the correct version.

On Mon, May 06, 2013 at 08:45:52AM +0800, Homer Hsing wrote:
> test cases for image3d_t
> 
> Signed-off-by: Homer Hsing <homer.xing at intel.com>
> ---
>  kernels/test_copy_image_3d.cl       | 11 ++++++++
>  kernels/test_fill_image_3d.cl       | 14 ++++++++++
>  kernels/test_fill_image_3d_2.cl     | 10 +++++++
>  utests/CMakeLists.txt               |  3 ++
>  utests/compiler_copy_image_3d.cpp   | 55 +++++++++++++++++++++++++++++++++++++
>  utests/compiler_fill_image_3d.cpp   | 44 +++++++++++++++++++++++++++++
>  utests/compiler_fill_image_3d_2.cpp | 42 ++++++++++++++++++++++++++++
>  7 files changed, 179 insertions(+)
>  create mode 100644 kernels/test_copy_image_3d.cl  create mode 100644 
> kernels/test_fill_image_3d.cl  create mode 100644 
> kernels/test_fill_image_3d_2.cl  create mode 100644 
> utests/compiler_copy_image_3d.cpp  create mode 100644 
> utests/compiler_fill_image_3d.cpp  create mode 100644 
> utests/compiler_fill_image_3d_2.cpp
> 
> diff --git a/kernels/test_copy_image_3d.cl 
> b/kernels/test_copy_image_3d.cl new file mode 100644 index 
> 0000000..766227a
> --- /dev/null
> +++ b/kernels/test_copy_image_3d.cl
> @@ -0,0 +1,11 @@
> +__kernel void
> +test_copy_image_3d(__read_only image3d_t src, __write_only image3d_t 
> +dst, sampler_t sampler) {
> +  int4 coord;
> +  int4 color;
> +  coord.x = (int)get_global_id(0);
> +  coord.y = (int)get_global_id(1);
> +  coord.z = 0;
> +  color = read_imagei(src, sampler, coord);
> +  write_imagei(dst, coord, color);
> +}
> diff --git a/kernels/test_fill_image_3d.cl 
> b/kernels/test_fill_image_3d.cl new file mode 100644 index 
> 0000000..0f0c6fd
> --- /dev/null
> +++ b/kernels/test_fill_image_3d.cl
> @@ -0,0 +1,14 @@
> +__kernel void
> +test_fill_image_3d(__write_only image3d_t dst, uint color) {
> +  int4 coord;
> +  int4 color4;
> +  color4.s0  = (color >> 24) & 0xFF;
> +  color4.s1  = (color >> 16) & 0xFF;
> +  color4.s2  = (color >> 8) & 0xFF;
> +  color4.s3  = color & 0xFF;
> +  coord.x = (int)get_global_id(0);
> +  coord.y = (int)get_global_id(1);
> +  coord.z = 0;
> +  write_imagei(dst, coord, color4);
> +}
> diff --git a/kernels/test_fill_image_3d_2.cl 
> b/kernels/test_fill_image_3d_2.cl new file mode 100644 index 
> 0000000..22b6452
> --- /dev/null
> +++ b/kernels/test_fill_image_3d_2.cl
> @@ -0,0 +1,10 @@
> +__kernel void
> +test_fill_image_3d_2(__write_only image3d_t dst) {
> +  int4 coord;
> +  int4 color4 = {0x12, 0x34, 0x56, 0x78};
> +  coord.x = (int)get_global_id(0);
> +  coord.y = (int)get_global_id(1);
> +  coord.z = 0;
> +  write_imagei(dst, coord, color4);
> +}
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index 
> a02d91f..1cecec8 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -23,10 +23,13 @@ set (utests_sources
>    compiler_convert_uchar_sat.cpp
>    compiler_copy_buffer.cpp
>    compiler_copy_image.cpp
> +  compiler_copy_image_3d.cpp
>    compiler_copy_buffer_row.cpp
>    compiler_fabs.cpp
>    compiler_fill_image.cpp
>    compiler_fill_image0.cpp
> +  compiler_fill_image_3d.cpp
> +  compiler_fill_image_3d_2.cpp
>    compiler_function_argument0.cpp
>    compiler_function_argument1.cpp
>    compiler_function_argument.cpp
> diff --git a/utests/compiler_copy_image_3d.cpp 
> b/utests/compiler_copy_image_3d.cpp
> new file mode 100644
> index 0000000..9816fe4
> --- /dev/null
> +++ b/utests/compiler_copy_image_3d.cpp
> @@ -0,0 +1,55 @@
> +#include "utest_helper.hpp"
> +
> +static void compiler_copy_image_3d(void) {
> +  const size_t w = 512;
> +  const size_t h = 512;
> +  const size_t depth = 1;
> +  cl_image_format format;
> +  cl_image_desc desc;
> +  cl_sampler sampler;
> +
> +  // Setup kernel and images
> +  OCL_CREATE_KERNEL("test_copy_image_3d");
> +  buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * w * h * depth);  
> + for (uint32_t k = 0; k < depth; k++)
> +    for (uint32_t j = 0; j < h; j++)
> +      for (uint32_t i = 0; i < w; i++)
> +        ((uint32_t*)buf_data[0])[k*w*h + j*w + i] = k*w*h + j*w + i;
> +
> +  format.image_channel_order = CL_RGBA;  
> + format.image_channel_data_type = CL_UNSIGNED_INT8;  desc.image_type 
> + = CL_MEM_OBJECT_IMAGE3D;  desc.image_width = w;  desc.image_height = 
> + h;  desc.image_depth = depth;  desc.image_row_pitch = 0;  
> + desc.image_slice_pitch = 0;  OCL_CREATE_IMAGE(buf[0], 
> + CL_MEM_COPY_HOST_PTR, &format, &desc, buf_data[0]);  
> + OCL_CREATE_IMAGE(buf[1], 0, &format, &desc, NULL);  
> + OCL_CREATE_SAMPLER(sampler, CL_ADDRESS_REPEAT, CL_FILTER_NEAREST);  
> + free(buf_data[0]);  buf_data[0] = 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(sampler), 
> + &sampler);  globals[0] = w;  globals[1] = h;  locals[0] = 16;  
> + locals[1] = 16;  OCL_NDRANGE(2);
> +
> +  // Check result
> +  OCL_MAP_BUFFER(0);
> +  OCL_MAP_BUFFER(1);
> +  for (uint32_t k = 0; k < depth; k++)
> +    for (uint32_t j = 0; j < h; ++j)
> +      for (uint32_t i = 0; i < w; i++)
> +        OCL_ASSERT(((uint32_t*)buf_data[0])[k*w*h + j*w + i] == 
> +((uint32_t*)buf_data[1])[k*w*h + j*w + i]);
> +  OCL_UNMAP_BUFFER(0);
> +  OCL_UNMAP_BUFFER(1);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_copy_image_3d);
> diff --git a/utests/compiler_fill_image_3d.cpp 
> b/utests/compiler_fill_image_3d.cpp
> new file mode 100644
> index 0000000..5d98b0b
> --- /dev/null
> +++ b/utests/compiler_fill_image_3d.cpp
> @@ -0,0 +1,44 @@
> +#include "utest_helper.hpp"
> +
> +static void compiler_fill_image_3d(void) {
> +  const size_t w = 512;
> +  const size_t h = 512;
> +  const size_t depth = 1;
> +  uint32_t color = 0x12345678;
> +  cl_image_format format;
> +  cl_image_desc desc;
> +
> +  format.image_channel_order = CL_RGBA;  
> + format.image_channel_data_type = CL_UNSIGNED_INT8;  desc.image_type 
> + = CL_MEM_OBJECT_IMAGE3D;  desc.image_width = w;  desc.image_height = 
> + h;  desc.image_depth = depth;  desc.image_row_pitch = 0;  
> + desc.image_slice_pitch = 0;
> +
> +  // Setup kernel and images
> +  OCL_CREATE_KERNEL("test_fill_image_3d");
> +
> +  OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
> +
> +  // Run the kernel
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);  OCL_SET_ARG(1, 
> + sizeof(color), &color);  globals[0] = w;  globals[1] = h;  locals[0] 
> + = 16;  locals[1] = 16;  OCL_NDRANGE(2);
> +
> +  // Check result
> +  OCL_MAP_BUFFER(0);
> +  for (uint32_t k = 0; k < depth; k++)
> +    for (uint32_t j = 0; j < h; ++j)
> +      for (uint32_t i = 0; i < w; i++)
> +        OCL_ASSERT(((uint32_t*)buf_data[0])[k*w*h + j*w + i] == 
> +0x78563412);
> +  OCL_UNMAP_BUFFER(0);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_3d);
> diff --git a/utests/compiler_fill_image_3d_2.cpp 
> b/utests/compiler_fill_image_3d_2.cpp
> new file mode 100644
> index 0000000..3b4edb9
> --- /dev/null
> +++ b/utests/compiler_fill_image_3d_2.cpp
> @@ -0,0 +1,42 @@
> +#include "utest_helper.hpp"
> +
> +static void compiler_fill_image_3d_2(void) {
> +  const size_t w = 512;
> +  const size_t h = 512;
> +  const size_t depth = 1;
> +  cl_image_format format;
> +  cl_image_desc desc;
> +
> +  format.image_channel_order = CL_RGBA;  
> + format.image_channel_data_type = CL_UNSIGNED_INT8;  desc.image_type 
> + = CL_MEM_OBJECT_IMAGE3D;  desc.image_width = w;  desc.image_height = 
> + h;  desc.image_depth = depth;  desc.image_row_pitch = 0;  
> + desc.image_slice_pitch = 0;
> +
> +  // Setup kernel and images
> +  OCL_CREATE_KERNEL("test_fill_image_3d_2");
> +
> +  OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL);
> +
> +  // Run the kernel
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);  globals[0] = w;  
> + globals[1] = h;  locals[0] = 16;  locals[1] = 16;  OCL_NDRANGE(2);
> +
> +  // Check result
> +  OCL_MAP_BUFFER(0);
> +  for (uint32_t k = 0; k < depth; k++)
> +    for (uint32_t j = 0; j < h; ++j)
> +      for (uint32_t i = 0; i < w; i++)
> +        OCL_ASSERT(((uint32_t*)buf_data[0])[k*w*h + j*w + i] == 
> +0x78563412);
> +  OCL_UNMAP_BUFFER(0);
> +}
> +
> +MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_3d_2);
> --
> 1.8.1.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list