This version LGTM, just pushed, thanks.
On Wed, Oct 22, 2014 at 03:51:06PM +0800, [email protected] wrote: > From: Junyan He <[email protected]> > > Signed-off-by: Junyan He <[email protected]> > --- > 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 > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
