[Beignet] [PATCH 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 | 44 +++++++++++++++++++++----------
 utests/CMakeLists.txt                 |  1 +
 utests/compiler_vector_load_store.cpp | 49 +++++++++++++++++++++++++++++++++--
 3 files changed, 78 insertions(+), 16 deletions(-)

diff --git a/kernels/compiler_vector_load_store.cl b/kernels/compiler_vector_load_store.cl
index 28fd93a..6a96662 100644
--- a/kernels/compiler_vector_load_store.cl
+++ b/kernels/compiler_vector_load_store.cl
@@ -1,18 +1,34 @@
 /* test OpenCL 1.1 Vector Data Load/Store Functions (section 6.11.7) */
-kernel void compiler_vector_load_store() {
-  float p[16], f;
+kernel void compiler_vector_load_store(__global float *p,
+                                       __global float *outf,
+                                       __global short *ps,
+                                       __global short *outs,
+                                       __global char *pchar,
+                                       __global char *outc) {
+  int x = get_global_id(0);
   float4 f4;
-  f4 = vload4(0, p);
-  vstore4(f4, 0, p);
-  
-  long x[16], l;
-  long16 l16;
-  l = vload16(0, x);
-  vstore16(l16, 0, x);
+  f4 = vload4(x, p);
+  f4 += (float4) {1.0f, 2.0f, 3.0f, 4.0f};
+  vstore4(f4, x, outf);
 
-  half h[16];
-  half4 h4;
-  f = vload_half(0, h);
-  f4 = vload_half4(0, h);
-  vstore_half(f, 0, h);
+  short4 s4;
+  s4 = vload4(x, ps);
+  s4 += (short4) {1, 2, 3, 4};
+  vstore4(s4, x, outs);
+
+  char4 c4;
+  c4 = vload4(x, pchar);
+  c4 += (char4) {1, 2, 3, 4};
+  vstore4(c4, x, outc);
+#if 0
+  long4 l4;
+  l4 = vload4(x, pl);
+  l4 += (long4) {1, 2, 3, 4};
+  vstore4(l4, x, outl);
+
+  half h4;
+  h4 = vload4(x, ph);
+  h4 += (long4) {1.0f, 2.0f, 3.0f, 4.0f};
+  vstore4(h4, x, outh);
+#endif
 }
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..17c4486 100644
--- a/utests/compiler_vector_load_store.cpp
+++ b/utests/compiler_vector_load_store.cpp
@@ -1,10 +1,55 @@
 #include "utest_helper.hpp"
 
-void compiler_vector_load_store(void)
+static void compiler_vector_load_store(void)
 {
+  const size_t n = 4 * 256;
+
+  // Setup kernel and buffers
   OCL_CREATE_KERNEL("compiler_vector_load_store");
+  buf_data[0] = (float*) malloc(sizeof(float) * n);
+  buf_data[1] = (short*) malloc(sizeof(short) * n);
+  buf_data[2] = (char*) malloc(sizeof(char) * n);
+  for (uint32_t i = 0; i < n; ++i)
+  {
+    ((float*)buf_data[0])[i] = i;
+    ((short*)buf_data[1])[i] = i & 0xFFFF;
+    ((char*)buf_data[2])[i] = i & 0xFF;
+  }
+  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);
+  OCL_CREATE_BUFFER(buf[2], CL_MEM_COPY_HOST_PTR, n * sizeof(short), buf_data[1]);
+  OCL_CREATE_BUFFER(buf[3], 0, n * sizeof(short), NULL);
+  OCL_CREATE_BUFFER(buf[4], CL_MEM_COPY_HOST_PTR, n * sizeof(char), buf_data[2]);
+  OCL_CREATE_BUFFER(buf[5], 0, n * sizeof(char), NULL);
+
+  // Run the kernel
+  for(uint32_t argID = 0; argID < 6; argID++)
+    OCL_SET_ARG(argID, sizeof(cl_mem), &buf[argID]);
+  globals[0] = n / 4;
+  locals[0] = 16;
+  OCL_NDRANGE(1);
+  free(buf_data[0]);
+  free(buf_data[1]);
+  free(buf_data[2]);
+  buf_data[0] = NULL;
+  buf_data[1] = NULL;
+  buf_data[2] = NULL;
+
+  // Check result
+  OCL_MAP_BUFFER(0);
+  OCL_MAP_BUFFER(1);
+  OCL_MAP_BUFFER(2);
+  OCL_MAP_BUFFER(3);
+  OCL_MAP_BUFFER(4);
+  OCL_MAP_BUFFER(5);
+  for (uint32_t i = 0; i < n; ++i)
+  {
+    int shift = ((i % 4) + 1);
+    OCL_ASSERT(((float*)buf_data[1])[i] == ((float*)buf_data[0])[i] + shift);
+    OCL_ASSERT(((short*)buf_data[3])[i] == ((short*)buf_data[2])[i] + shift);
+    OCL_ASSERT(((char*)buf_data[5])[i] == (char)(((char*)buf_data[4])[i] + (char)shift));
+  }
 }
 
 MAKE_UTEST_FROM_FUNCTION(compiler_vector_load_store);
 
-
-- 
1.7.11.7