[Beignet] [PATCH v2 2/2] utests: test vector load and store.
Zhigang Gong
zhigang.gong at linux.intel.com
Fri May 24 04:07:17 PDT 2013
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
More information about the Beignet
mailing list