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

[Zhigang Gong]

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