[Beignet] [PATCH 3/3] Add a test case that trigger vstore3 bug.

[Chuanbo Weng]

When there is no workgroup, the vstore3 will store incorrect
value to buffer.

Signed-off-by: Chuanbo Weng <chuanbo.weng at intel.com>
---
 kernels/compiler_vstore3_no_workgroup.cl |   11 +++++++++
 utests/CMakeLists.txt                    |    1 +
 utests/compiler_vstore3_no_workgroup.cpp |   39 ++++++++++++++++++++++++++++++
 3 files changed, 51 insertions(+)
 create mode 100644 kernels/compiler_vstore3_no_workgroup.cl
 create mode 100644 utests/compiler_vstore3_no_workgroup.cpp

diff --git a/kernels/compiler_vstore3_no_workgroup.cl b/kernels/compiler_vstore3_no_workgroup.cl
new file mode 100644
index 0000000..aa641ae
--- /dev/null
+++ b/kernels/compiler_vstore3_no_workgroup.cl
@@ -0,0 +1,11 @@
+__kernel
+void compiler_vstore3_no_workgroup(__global uint *vectorArray){
+  int col = get_global_id(0);
+
+  uint3 val3;
+  val3.x = col*3 + 0;
+  val3.y = col*3 + 1;
+  val3.z = col*3 + 2;
+  size_t thumb_offset = col;
+  vstore3(val3, thumb_offset, vectorArray);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index 98db273..b0ac53a 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -136,6 +136,7 @@ set (utests_sources
   compiler_long_mult.cpp
   compiler_long_cmp.cpp
   compiler_bool_cross_basic_block.cpp
+  compiler_vstore3_no_workgroup.cpp
   utest_assert.cpp
   utest.cpp
   utest_file_map.cpp
diff --git a/utests/compiler_vstore3_no_workgroup.cpp b/utests/compiler_vstore3_no_workgroup.cpp
new file mode 100644
index 0000000..96491cd
--- /dev/null
+++ b/utests/compiler_vstore3_no_workgroup.cpp
@@ -0,0 +1,39 @@
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, unsigned int *dst) {
+  for(unsigned int j = 0; j < 3; j++){
+    *(dst + global_id*3 + j) = global_id*3 + j;
+  }
+}
+
+void compiler_vstore3_no_workgroup(void){
+
+  const size_t n = 16;
+  unsigned int cpu_dst[16*3];
+	
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL("compiler_vstore3_no_workgroup");
+  OCL_CREATE_BUFFER(buf[0], 0, n*3*sizeof(unsigned int), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  globals[0] = n;
+
+  // Run the kernel on GPU
+  OCL_CALL (clEnqueueNDRangeKernel, queue, kernel, 1, NULL, globals, NULL, 0, NULL, NULL);
+
+  // Run on CPU
+  for (int32_t i = 0; i < (int32_t) n; ++i)
+    cpu(i, cpu_dst);
+
+  // Compare
+  OCL_MAP_BUFFER(0);
+  unsigned int *px = (unsigned int *)buf_data[0];
+  for (int32_t i = 0; i < (int32_t) n; ++i)
+    for (int32_t j = 0; j < 3; ++j){
+      OCL_ASSERT(*(px+i*3+j) == *(cpu_dst+i*3+j));
+    }
+  OCL_UNMAP_BUFFER(0);
+
+}
+
+
+MAKE_UTEST_FROM_FUNCTION(compiler_vstore3_no_workgroup)
-- 
1.7.9.5