LGTM, will push latter.
On Thu, May 15, 2014 at 04:43:37PM +0800, [email protected] wrote: > From: Junyan He <[email protected]> > > Signed-off-by: Junyan He <[email protected]> > --- > kernels/test_copy_image_1d.cl | 9 +++++++ > kernels/test_fill_image_1d.cl | 8 ++++++ > utests/CMakeLists.txt | 2 ++ > utests/compiler_copy_image_1d.cpp | 52 > +++++++++++++++++++++++++++++++++++++++ > utests/compiler_fill_image_1d.cpp | 44 +++++++++++++++++++++++++++++++++ > 5 files changed, 115 insertions(+) > create mode 100644 kernels/test_copy_image_1d.cl > create mode 100644 kernels/test_fill_image_1d.cl > create mode 100644 utests/compiler_copy_image_1d.cpp > create mode 100644 utests/compiler_fill_image_1d.cpp > > diff --git a/kernels/test_copy_image_1d.cl b/kernels/test_copy_image_1d.cl > new file mode 100644 > index 0000000..88428bb > --- /dev/null > +++ b/kernels/test_copy_image_1d.cl > @@ -0,0 +1,9 @@ > +__kernel void > +test_copy_image_1d(__read_only image1d_t src, __write_only image1d_t dst, > sampler_t sampler) > +{ > + int coord; > + int4 color; > + coord = (int)get_global_id(0); > + color = read_imagei(src, sampler, coord); > + write_imagei(dst, coord, color); > +} > diff --git a/kernels/test_fill_image_1d.cl b/kernels/test_fill_image_1d.cl > new file mode 100644 > index 0000000..db922af > --- /dev/null > +++ b/kernels/test_fill_image_1d.cl > @@ -0,0 +1,8 @@ > +__kernel void > +test_fill_image_1d(__write_only image1d_t dst) > +{ > + int coord; > + coord = (int)get_global_id(0); > + uint4 color4 = {0, 1, 2 ,3}; > + write_imageui(dst, coord, color4); > +} > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt > index 415dcb6..91e4ffb 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -45,6 +45,7 @@ set (utests_sources > compiler_convert_uchar_sat.cpp > compiler_copy_buffer.cpp > compiler_copy_image.cpp > + compiler_copy_image_1d.cpp > compiler_copy_image_3d.cpp > compiler_copy_buffer_row.cpp > compiler_degrees.cpp > @@ -54,6 +55,7 @@ set (utests_sources > compiler_abs_diff.cpp > compiler_fill_image.cpp > compiler_fill_image0.cpp > + compiler_fill_image_1d.cpp > compiler_fill_image_3d.cpp > compiler_fill_image_3d_2.cpp > compiler_function_argument0.cpp > diff --git a/utests/compiler_copy_image_1d.cpp > b/utests/compiler_copy_image_1d.cpp > new file mode 100644 > index 0000000..5af6a77 > --- /dev/null > +++ b/utests/compiler_copy_image_1d.cpp > @@ -0,0 +1,52 @@ > +#include <string.h> > +#include "utest_helper.hpp" > + > +static void compiler_copy_image_1d(void) > +{ > + const size_t w = 512; > + cl_image_format format; > + cl_image_desc desc; > + cl_sampler sampler; > + > + memset(&desc, 0x0, sizeof(cl_image_desc)); > + memset(&format, 0x0, sizeof(cl_image_format)); > + > + // Setup kernel and images > + OCL_CREATE_KERNEL("test_copy_image_1d"); > + buf_data[0] = (uint32_t*) malloc(sizeof(uint32_t) * w); > + for (uint32_t i = 0; i < w; i++) > + ((uint32_t*)buf_data[0])[i] = i; > + > + format.image_channel_order = CL_RGBA; > + format.image_channel_data_type = CL_UNSIGNED_INT8; > + desc.image_type = CL_MEM_OBJECT_IMAGE1D; > + desc.image_width = w; > + desc.image_row_pitch = w * sizeof(uint32_t); > + OCL_CREATE_IMAGE(buf[0], CL_MEM_COPY_HOST_PTR, &format, &desc, > buf_data[0]); > + > + desc.image_row_pitch = 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; > + locals[0] = 16; > + OCL_NDRANGE(1); > + > + // Check result > + OCL_MAP_BUFFER(0); > + OCL_MAP_BUFFER(1); > + for (uint32_t i = 0; i < w; i++) { > + //printf (" %x", ((uint32_t*)buf_data[1])[i]); > + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == ((uint32_t*)buf_data[1])[i]); > + } > + OCL_UNMAP_BUFFER(0); > + OCL_UNMAP_BUFFER(1); > +} > + > +MAKE_UTEST_FROM_FUNCTION(compiler_copy_image_1d); > diff --git a/utests/compiler_fill_image_1d.cpp > b/utests/compiler_fill_image_1d.cpp > new file mode 100644 > index 0000000..c1011ba > --- /dev/null > +++ b/utests/compiler_fill_image_1d.cpp > @@ -0,0 +1,44 @@ > +#include <string.h> > +#include "utest_helper.hpp" > + > +static void compiler_fill_image_1d(void) > +{ > + const size_t w = 2048; > + cl_image_format format; > + cl_image_desc desc; > + > + 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_IMAGE1D; > + desc.image_width = w; > + desc.image_row_pitch = 0; > + > + // Setup kernel and images > + OCL_CREATE_KERNEL("test_fill_image_1d"); > + > + OCL_CREATE_IMAGE(buf[0], 0, &format, &desc, NULL); > + > + // Run the kernel > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + globals[0] = w/2; > + locals[0] = 16; > + OCL_NDRANGE(1); > + > + // Check result > + OCL_MAP_BUFFER_GTT(0); > + //printf("------ The image result is: -------\n"); > + for (uint32_t i = 0; i < w/2; i++) { > + //printf(" %2x", ((uint32_t *)buf_data[0])[i]); > + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 0x03020100); > + } > + for (uint32_t i = w/2; i < w; i++) { > + //printf(" %2x", ((uint32_t *)buf_data[0])[i]); > + OCL_ASSERT(((uint32_t*)buf_data[0])[i] == 0); > + } > + OCL_UNMAP_BUFFER_GTT(0); > +} > + > +MAKE_UTEST_FROM_FUNCTION(compiler_fill_image_1d); > -- > 1.8.3.2 > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
