[Beignet] [PATCH 2/2] utests: Add a unit test for non-aligned group size.

[Zhigang Gong]

On Thu, Aug 22, 2013 at 04:52:05PM +0800, Ruiling Song wrote:
> To hit prediction logic.
> 
> Signed-off-by: Ruiling Song <ruiling.song at intel.com>
> ---
>  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
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet