[Beignet] [PATCH 5/5] test cases for image3d_t
Zhigang Gong
zhigang.gong at linux.intel.com
Tue May 7 22:04:38 PDT 2013
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
-------------- next part --------------
A non-text attachment was scrubbed...
Name: 0001-GBE-fixed-the-hard-coded-implementation-for-sampler-.patch
Type: application/octet-stream
Size: 4655 bytes
Desc: not available
URL: <http://lists.freedesktop.org/archives/beignet/attachments/20130508/ad62e8ca/attachment-0001.obj>
More information about the Beignet
mailing list