Sorry for the inconvenience, I have sent a patch to fix it. Please have a try.
-----Original Message----- From: Zhigang Gong [mailto:[email protected]] Sent: Friday, August 30, 2013 4:15 PM To: Song, Ruiling Cc: [email protected] Subject: Re: [Beignet] [PATCH 2/2] utests: Add a unit test for non-aligned group size. On Thu, Aug 22, 2013 at 04:52:05PM +0800, Ruiling Song wrote: > To hit prediction logic. > > Signed-off-by: Ruiling Song <[email protected]> > --- > kernels/compiler_group_size.cl | 17 +++++++++++++ > utests/compiler_group_size.cpp | 52 > ++++++++++++++++++++++++++++++++++++++++ > 2 files changed, 69 insertions(+) > > diff --git a/kernels/compiler_group_size.cl > b/kernels/compiler_group_size.cl index 9dba236..0fe88c1 100644 > --- a/kernels/compiler_group_size.cl > +++ b/kernels/compiler_group_size.cl > @@ -10,3 +10,20 @@ compiler_group_size(__global unsigned int *dst) > dst[idz*size_x*size_y + idy*size_x + idx] = idz*size_x*size_y + > idy*size_x +idx; } > > +struct xyz{ > + unsigned b:16; > + unsigned e:16; > + unsigned o; > +}; Per opencl spec, it doesn't support bitfileds. I met the following errors with LLVM3.3 [CL_RGBx CL_UNORM_INT_101010] compiler_group_size1: /tmp/file4DjvoO.cl:14:12: error: bitfields are not supported in OpenCL unsigned b:16; ^ /tmp/file4DjvoO.cl:15:12: error: bitfields are not supported in OpenCL unsigned e:16; ^ 2 errors generated. compiler_group_size1() [FAILED] Error: error calling clBuildProgram with errorCL_INVALID_PROGRAM I did not noticed that error and already pushed this patch to the master tree. Could you make another patch to fix this issue? Thanks. > + > +__kernel void > +compiler_group_size4(__global struct xyz *src, __global unsigned int > +*dst, unsigned int num, unsigned int c) { > + uint idx = (uint)get_global_id(0); > + if(idx>=num) > + return; > + struct xyz td = src[idx]; > + for(unsigned x = td.b;x<=td.e;x++) > + dst[td.o+x] = c; > +} > + > diff --git a/utests/compiler_group_size.cpp > b/utests/compiler_group_size.cpp index 6d59aed..02544b2 100644 > --- a/utests/compiler_group_size.cpp > +++ b/utests/compiler_group_size.cpp > @@ -1,4 +1,11 @@ > #include "utest_helper.hpp" > +#include <string.h> > + > +struct xyz{ > + unsigned b:16; > + unsigned e:16; > + unsigned o; > +}; > > void compiler_group_size1(void) > { > @@ -80,7 +87,52 @@ void compiler_group_size3(void) > OCL_UNMAP_BUFFER(0); > } > } > + > +void compiler_group_size4(void) > +{ > + const size_t n = 16; > + uint32_t color = 2; > + uint32_t num = 1; > + int group_size[] = {1}; > + // Setup kernel and buffers > + OCL_CREATE_KERNEL_FROM_FILE("compiler_group_size", > +"compiler_group_size4"); > + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(struct xyz), NULL); > + OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL); > + > + for(uint32_t i = 0; i < num; i++) { > + // Run the kernel > + OCL_MAP_BUFFER(0); > + ((struct xyz*)buf_data[0])[0].b = 0; > + ((struct xyz*)buf_data[0])[0].e = 2; > + ((struct xyz*)buf_data[0])[0].o = 0; > + OCL_UNMAP_BUFFER(0); > + > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); > + OCL_SET_ARG(2, sizeof(cl_int), &group_size[i]); > + OCL_SET_ARG(3, sizeof(cl_int), &color); > + > + globals[0] = group_size[i]; > + locals[0] = group_size[i]; > + OCL_NDRANGE(1); > + OCL_MAP_BUFFER(1); > + > + // Check results > + for (uint32_t j = 0; j < n; ++j) { > +// std::cout <<((uint32_t*)buf_data[1])[j] << " "; > + if(j >= i && j <= i+2) { > + OCL_ASSERT(((uint32_t*)buf_data[1])[j] == color); > + } else { > + OCL_ASSERT(((uint32_t*)buf_data[1])[j] == 0); > + } > + > + } > + memset(((uint32_t*)buf_data[1]), 0x0, sizeof(int)*n); > + OCL_UNMAP_BUFFER(1); > + } > +} > MAKE_UTEST_FROM_FUNCTION(compiler_group_size1); > MAKE_UTEST_FROM_FUNCTION(compiler_group_size2); > MAKE_UTEST_FROM_FUNCTION(compiler_group_size3); > +MAKE_UTEST_FROM_FUNCTION(compiler_group_size4); > > -- > 1.7.9.5 > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
