[Beignet] [PATCH 5/5] test cases for image3d_t
Xing, Homer
homer.xing at intel.com
Tue May 7 22:19:51 PDT 2013
That patch has been tried. It does not break any test case.
Homer
-----Original Message-----
From: beignet-bounces+homer.xing=intel.com at lists.freedesktop.org [mailto:beignet-bounces+homer.xing=intel.com at lists.freedesktop.org] On Behalf Of Xing, Homer
Sent: Wednesday, May 08, 2013 1:16 PM
To: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH 5/5] test cases for image3d_t
I am glad to see a potential bug has been fixed by you.
I have reviewed your patch. You may delete "printf("sampler dst num %d src num %d \n", insn.getDstNum(), insn.getSrcNum());".
I will test it and tell you the result.
Homer
-----Original Message-----
From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com]
Sent: Wednesday, May 08, 2013 1:05 PM
To: Xing, Homer; beignet at lists.freedesktop.org
Subject: RE: [Beignet] [PATCH 5/5] test cases for image3d_t
This patchset is based on the previous image2d implementation. I just reviewed the patchset and found you didn't change anything in gen_insn_selection.cpp. As there are some hard coded impelemtnations in both the Sampler and TypedWrite instructions' implementations in the insn selection stage. You have to modify it to keep consistent with the image3d implementation, otherwise there are some fatal errors including using incorrect registers in the code gen stage and some buffer overflows. You must be very lucky guy that it can work well in your platform. In my platform, it always has the following three random results: :) segfault/assert/incorrect value.
The attached file is a patch to gen_insn_selection.cpp which is to fix those hard coded implementation please review it and have a try. This patch should be applied before this image3d_t ptachset and after that the image3d patchset works well on my platform.
Thanks.
> -----Original Message-----
> From:
> beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org
> [mailto:beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.
> org] On Behalf Of Xing, Homer
> Sent: Wednesday, May 08, 2013 8:25 AM
> To: beignet at lists.freedesktop.org
> Subject: Re: [Beignet] [PATCH 5/5] test cases for image3d_t
>
> 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
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list