[Beignet] [PATCH v2 2/2] utests: test vector load and store.

Yang, Rong R rong.r.yang at intel.com
Tue May 28 22:24:15 PDT 2013


LGTM.

-----Original Message-----
From: beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org [mailto:beignet-bounces+rong.r.yang=intel.com at lists.freedesktop.org] On Behalf Of Zhigang Gong
Sent: Friday, May 24, 2013 7:07 PM
To: beignet at lists.freedesktop.org
Subject: Re: [Beignet] [PATCH v2 2/2] utests: test vector load and store.

It seems that we have bugs on the 3 elements vector, so disable it right now.
Will enable it after we identify and fix that bug.

> -----Original Message-----
> From: 
> beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org
>
[mailto:beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org]
> On Behalf Of Zhigang Gong
> Sent: Friday, May 24, 2013 7:06 PM
> To: beignet at lists.freedesktop.org
> Cc: Zhigang Gong
> Subject: [Beignet] [PATCH v2 2/2] utests: test vector load and store.
> 
> Add float4/short4/char4 test case.
> 
> Signed-off-by: Zhigang Gong <zhigang.gong at linux.intel.com>
> ---
>  kernels/compiler_vector_load_store.cl | 52
> +++++++++++++++++++++++---------
>  utests/CMakeLists.txt                 |  1 +
>  utests/compiler_vector_load_store.cpp | 57
> ++++++++++++++++++++++++++++++++---
>  3 files changed, 91 insertions(+), 19 deletions(-)
> 
> diff --git a/kernels/compiler_vector_load_store.cl
> b/kernels/compiler_vector_load_store.cl
> index 28fd93a..b362412 100644
> --- a/kernels/compiler_vector_load_store.cl
> +++ b/kernels/compiler_vector_load_store.cl
> @@ -1,18 +1,40 @@
>  /* test OpenCL 1.1 Vector Data Load/Store Functions (section 6.11.7) 
> */ -kernel void compiler_vector_load_store() {
> -  float p[16], f;
> -  float4 f4;
> -  f4 = vload4(0, p);
> -  vstore4(f4, 0, p);
> -
> -  long x[16], l;
> -  long16 l16;
> -  l = vload16(0, x);
> -  vstore16(l16, 0, x);
> 
> -  half h[16];
> -  half4 h4;
> -  f = vload_half(0, h);
> -  f4 = vload_half4(0, h);
> -  vstore_half(f, 0, h);
> +#define OFFSET2(type)  (type ##2) {(type)1, (type)2} #define
> +OFFSET3(type)  (type ##3) {(type)1, (type)2, (type)3} #define
> +OFFSET4(type)  (type ##4) {(type)1, (type)2, (type)3, (type)4} 
> +#define
> +OFFSET8(type)  (type ##8) {(type)1, (type)2, (type)3, (type)4, 
> +(type)5, (type)6, (type)7, (type)8} #define OFFSET16(type)  (type 
> +##16) {(type)1, (type)2, (type)3, (type)4, (type)5, (type)6, (type)7, 
> +(type)8, (type)9, (type)10, (type)11, (type)12, (type)13, (type)14, 
> +(type)15, (type)16}
> +
> +#define  TEST_TYPE(type, n) \
> +__kernel void test_##type ##n(__global type *pin, \
> +                            __global type *pout)  \ {\
> +  int x = get_global_id(0); \
> +  type ##n value; \
> +  value = vload ##n(x, pin); \
> +  value += OFFSET ##n(type); \
> +  vstore ##n(value, x, pout); \
>  }
> +
> +#define TEST_ALL_TYPE(n) \
> +  TEST_TYPE(char,n) \
> +  TEST_TYPE(uchar,n) \
> +  TEST_TYPE(short,n) \
> +  TEST_TYPE(ushort,n) \
> +  TEST_TYPE(int,n) \
> +  TEST_TYPE(uint,n) \
> +  TEST_TYPE(float,n)
> +
> +#if 0
> +  TEST_TYPE(double,n)
> +  TEST_TYPE(long,n)
> +  TEST_TYPE(ulong,n)
> +  TEST_TYPE(half,n)
> +#endif
> +
> +TEST_ALL_TYPE(2)
> +//TEST_ALL_TYPE(3)
> +TEST_ALL_TYPE(4)
> +TEST_ALL_TYPE(8)
> +TEST_ALL_TYPE(16)
> diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index
> 63c873d..268cc51 100644
> --- a/utests/CMakeLists.txt
> +++ b/utests/CMakeLists.txt
> @@ -76,6 +76,7 @@ set (utests_sources
>    compiler_volatile.cpp
>    compiler_copy_image1.cpp
>    compiler_get_image_info.cpp
> +  compiler_vector_load_store.cpp
>    runtime_createcontext.cpp
>    utest_assert.cpp
>    utest.cpp
> diff --git a/utests/compiler_vector_load_store.cpp
> b/utests/compiler_vector_load_store.cpp
> index 96fcfa9..76c12a1 100644
> --- a/utests/compiler_vector_load_store.cpp
> +++ b/utests/compiler_vector_load_store.cpp
> @@ -1,10 +1,59 @@
>  #include "utest_helper.hpp"
> -
> -void compiler_vector_load_store(void)
> +template<typename T>
> +static void compiler_vector_load_store(int elemNum, const char
> +*kernelName)
>  {
> -  OCL_CREATE_KERNEL("compiler_vector_load_store");
> +  const size_t n = elemNum * 256;
> +
> +  // Setup kernel and buffers
> +  OCL_CREATE_KERNEL_FROM_FILE("compiler_vector_load_store",
> + kernelName);  buf_data[0] = (T*) malloc(sizeof(T) * n);  for 
> + (uint32_t i = 0; i < n; ++i)
> +    ((T*)buf_data[0])[i] = i;
> +  OCL_CREATE_BUFFER(buf[0], CL_MEM_COPY_HOST_PTR, n * sizeof(float), 
> + buf_data[0]);  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), 
> + NULL); 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]);  globals[0] = n / elemNum;  locals[0] = 
> + 16; OCL_NDRANGE(1);
> +
> +  // Check result
> +  OCL_MAP_BUFFER(0);
> +  OCL_MAP_BUFFER(1);
> +  for (uint32_t i = 0; i < n; ++i)
> +  {
> +    int shift = ((i % elemNum) + 1);
> +    OCL_ASSERT(((T*)buf_data[1])[i] == (T)(((T*)buf_data[0])[i] + 
> + shift));  }  OCL_UNMAP_BUFFER(0);  OCL_UNMAP_BUFFER(1);
>  }
> 
> -MAKE_UTEST_FROM_FUNCTION(compiler_vector_load_store);
> +#define compiler_vector_load_store(type, n, kernel_type) \ static 
> +void compiler_vector_ ##kernel_type ##n ##_load_store(void)\ {\
> +  compiler_vector_load_store<type>(n, "test_" #kernel_type #n);\ }\ 
> +MAKE_UTEST_FROM_FUNCTION(compiler_vector_ ## kernel_type ##n 
> +##_load_store);
> 
> +#define test_all_vector(type, kernel_type) \
> +  compiler_vector_load_store(type, 2, kernel_type) \
> +  /*compiler_vector_load_store(type, 3, kernel_type)*/ \
> +  compiler_vector_load_store(type, 4, kernel_type) \
> +  compiler_vector_load_store(type, 8, kernel_type) \
> +  compiler_vector_load_store(type, 16, kernel_type)
> 
> +test_all_vector(int8_t, char)
> +test_all_vector(uint8_t, uchar)
> +test_all_vector(int16_t, short)
> +test_all_vector(uint16_t, ushort)
> +test_all_vector(int32_t, int)
> +test_all_vector(uint32_t, uint)
> +test_all_vector(float, float)
> +//test_all_vector(double, double)
> +//test_all_vector(int64_t, long)
> +//test_all_vector(uint64_t, ulong)
> --
> 1.7.11.7
> 
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet

_______________________________________________
Beignet mailing list
Beignet at lists.freedesktop.org
http://lists.freedesktop.org/mailman/listinfo/beignet


More information about the Beignet mailing list