[Beignet] [PATCH] Add more type for async copy test case.

[Zhigang Gong]

LGTM, will push latter. Thanks.

On Mon, Oct 21, 2013 at 03:47:56PM +0800, Yang Rong wrote:
> 
> Signed-off-by: Yang Rong <rong.r.yang at intel.com>
> ---
>  kernels/compiler_async_copy.cl | 38 +++++++++++--------
>  utests/compiler_async_copy.cpp | 86 +++++++++++++++++++++++++-----------------
>  2 files changed, 74 insertions(+), 50 deletions(-)
> 
> diff --git a/kernels/compiler_async_copy.cl b/kernels/compiler_async_copy.cl
> index a2432a4..06ec8e7 100644
> --- a/kernels/compiler_async_copy.cl
> +++ b/kernels/compiler_async_copy.cl
> @@ -1,16 +1,24 @@
> -__kernel void
> -compiler_async_copy(__global int2 *dst, __global int2 *src, __local int2 *localBuffer, int copiesPerWorkItem)
> -{
> -  event_t event;
> -  int copiesPerWorkgroup = copiesPerWorkItem * get_local_size(0);
> -  int i;
> -  event = async_work_group_copy((__local int2*)localBuffer, (__global const int2*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, (event_t)0 );
> -  wait_group_events( 1, &event );
> -
> -  for(i=0; i<copiesPerWorkItem; i++)
> -    localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] + (int2)(3, 3);
> -  barrier(CLK_LOCAL_MEM_FENCE);
> -
> -  event = async_work_group_copy((__global int2*)(dst+copiesPerWorkgroup*get_group_id(0)), (__local const int2*)localBuffer, (size_t)copiesPerWorkgroup, (event_t)0 );
> -  wait_group_events( 1, &event );
> +#define DEF(TYPE) \
> +kernel void \
> +compiler_async_copy_##TYPE(__global TYPE *dst, __global TYPE *src, __local TYPE *localBuffer, int copiesPerWorkItem) \
> +{ \
> +  event_t event; \
> +  int copiesPerWorkgroup = copiesPerWorkItem * get_local_size(0); \
> +  int i; \
> +  event = async_work_group_copy((__local TYPE*)localBuffer, (__global const TYPE*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, (event_t)0 ); \
> +  wait_group_events( 1, &event ); \
> +\
> +  event = async_work_group_copy((__global TYPE*)(dst+copiesPerWorkgroup*get_group_id(0)), (__local const TYPE*)localBuffer, (size_t)copiesPerWorkgroup, (event_t)0 ); \
> +  wait_group_events( 1, &event ); \
>  }
> +
> +DEF(char2);
> +DEF(uchar2);
> +DEF(short2);
> +DEF(ushort2);
> +DEF(int2);
> +DEF(uint2);
> +DEF(long2);
> +DEF(ulong2);
> +DEF(float2);
> +DEF(double2);
> diff --git a/utests/compiler_async_copy.cpp b/utests/compiler_async_copy.cpp
> index 9384f85..7951ff7 100644
> --- a/utests/compiler_async_copy.cpp
> +++ b/utests/compiler_async_copy.cpp
> @@ -1,39 +1,55 @@
>  #include "utest_helper.hpp"
> +#include <stdint.h>
>  
> -static void compiler_async_copy(void)
> -{
> -  const size_t n = 1024;
> -  const size_t local_size = 32;
> -  const int copiesPerWorkItem = 5;
> +typedef unsigned char uchar;
> +typedef unsigned short ushort;
>  
> -  // Setup kernel and buffers
> -  OCL_CREATE_KERNEL("compiler_async_copy");
> -  OCL_CREATE_BUFFER(buf[0], 0, n * copiesPerWorkItem * sizeof(int) * 2, NULL);
> -  OCL_CREATE_BUFFER(buf[1], 0, n * copiesPerWorkItem * sizeof(int) * 2, NULL);
> -  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
> -  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
> -  OCL_SET_ARG(2, local_size*copiesPerWorkItem*sizeof(int)*2, NULL);
> -  OCL_SET_ARG(3, sizeof(int), &copiesPerWorkItem);
> +#define DEF(TYPE, KER_TYPE, VEC_SIZE) \
> +static void compiler_async_copy_##KER_TYPE##VEC_SIZE(void) \
> +{ \
> +  const size_t n = 1024; \
> +  const size_t local_size = 32; \
> +  const int copiesPerWorkItem = 5; \
> +\
> +  /* Setup kernel and buffers */\
> +  OCL_CREATE_KERNEL_FROM_FILE("compiler_async_copy", "compiler_async_copy_" # KER_TYPE # VEC_SIZE); \
> +  OCL_CREATE_BUFFER(buf[0], 0, n * copiesPerWorkItem * sizeof(TYPE) * VEC_SIZE, NULL); \
> +  OCL_CREATE_BUFFER(buf[1], 0, n * copiesPerWorkItem * sizeof(TYPE) * VEC_SIZE, NULL); \
> +  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); \
> +  OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]); \
> +  OCL_SET_ARG(2, local_size*copiesPerWorkItem*sizeof(TYPE)*VEC_SIZE, NULL); \
> +  OCL_SET_ARG(3, sizeof(int), &copiesPerWorkItem); \
> +\
> +  OCL_MAP_BUFFER(1); \
> +  for (uint32_t i = 0; i < n * copiesPerWorkItem * VEC_SIZE; ++i) \
> +      ((TYPE*)buf_data[1])[i] = rand(); \
> +  OCL_UNMAP_BUFFER(1); \
> +\
> +  /* Run the kernel */\
> +  globals[0] = n; \
> +  locals[0] = local_size; \
> +  OCL_NDRANGE(1); \
> +  OCL_MAP_BUFFER(0); \
> +  OCL_MAP_BUFFER(1); \
> +\
> +  /* Check results */\
> +  TYPE *dst = (TYPE*)buf_data[0]; \
> +  TYPE *src = (TYPE*)buf_data[1]; \
> +  for (uint32_t i = 0; i < n * copiesPerWorkItem * VEC_SIZE; i++) \
> +    OCL_ASSERT(dst[i] == src[i]); \
> +  OCL_UNMAP_BUFFER(0); \
> +  OCL_UNMAP_BUFFER(1); \
> +} \
> +\
> +MAKE_UTEST_FROM_FUNCTION(compiler_async_copy_##KER_TYPE##VEC_SIZE);
>  
> -  OCL_MAP_BUFFER(1);
> -  for (uint32_t i = 0; i < n * copiesPerWorkItem * 2; ++i)
> -      ((int*)buf_data[1])[i] = rand();
> -  OCL_UNMAP_BUFFER(1);
> -
> -  // Run the kernel
> -  globals[0] = n;
> -  locals[0] = local_size;
> -  OCL_NDRANGE(1);
> -  OCL_MAP_BUFFER(0);
> -  OCL_MAP_BUFFER(1);
> -
> -  // Check results
> -  int *dst = (int*)buf_data[0];
> -  int *src = (int*)buf_data[1];
> -  for (uint32_t i = 0; i < n * copiesPerWorkItem * 2; i++)
> -    OCL_ASSERT(dst[i] == src[i] + 3);
> -  OCL_UNMAP_BUFFER(0);
> -  OCL_UNMAP_BUFFER(1);
> -}
> -
> -MAKE_UTEST_FROM_FUNCTION(compiler_async_copy);
> +DEF(char, char, 2);
> +DEF(uchar, uchar, 2);
> +DEF(short, short, 2);
> +DEF(ushort, ushort, 2);
> +DEF(int, int, 2);
> +DEF(uint, uint, 2);
> +DEF(int64_t, long, 2);
> +DEF(uint64_t, ulong, 2);
> +DEF(float, float, 2);
> +DEF(double, double, 2);
> -- 
> 1.8.1.2
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet