utests: Add a unit test for non-aligned group size.

To hit prediction logic.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

diff --git a/kernels/compiler_group_size.cl b/kernels/compiler_group_size.cl
index 9dba236..0fe88c1 100644 (file)
--- a/kernels/compiler_group_size.cl
+++ b/kernels/compiler_group_size.cl
@@ -10,3 +10,20 @@ compiler_group_size(__global unsigned int *dst)
   dst[idz*size_x*size_y + idy*size_x + idx] = idz*size_x*size_y + idy*size_x +idx;
 }
 
+struct xyz{
+  unsigned b:16;
+  unsigned e:16;
+  unsigned o;
+};
+
+__kernel void
+compiler_group_size4(__global struct xyz *src, __global unsigned int *dst, unsigned int num, unsigned int c)
+{
+  uint idx = (uint)get_global_id(0);
+  if(idx>=num)
+    return;
+  struct xyz td = src[idx];
+  for(unsigned x = td.b;x<=td.e;x++)
+    dst[td.o+x] = c;
+}
+

diff --git a/utests/compiler_group_size.cpp b/utests/compiler_group_size.cpp
index 6d59aed..02544b2 100644 (file)
--- a/utests/compiler_group_size.cpp
+++ b/utests/compiler_group_size.cpp
@@ -1,4 +1,11 @@
 #include "utest_helper.hpp"
+#include <string.h>
+
+struct xyz{
+  unsigned b:16;
+  unsigned e:16;
+  unsigned o;
+};
 
 void compiler_group_size1(void)
 {
@@ -80,7 +87,52 @@ void compiler_group_size3(void)
     OCL_UNMAP_BUFFER(0);
   }
 }
+
+void compiler_group_size4(void)
+{
+  const size_t n = 16;
+  uint32_t color = 2;
+  uint32_t num = 1;
+  int group_size[] = {1};
+  // Setup kernel and buffers
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_group_size", "compiler_group_size4");
+  OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(struct xyz), NULL);
+  OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(uint32_t), NULL);
+
+  for(uint32_t i = 0; i < num; i++) {
+    // Run the kernel
+    OCL_MAP_BUFFER(0);
+    ((struct xyz*)buf_data[0])[0].b = 0;
+    ((struct xyz*)buf_data[0])[0].e = 2;
+    ((struct xyz*)buf_data[0])[0].o = 0;
+    OCL_UNMAP_BUFFER(0);
+
+    OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+    OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+    OCL_SET_ARG(2, sizeof(cl_int), &group_size[i]);
+    OCL_SET_ARG(3, sizeof(cl_int), &color);
+
+    globals[0] = group_size[i];
+    locals[0] = group_size[i];
+    OCL_NDRANGE(1);
+    OCL_MAP_BUFFER(1);
+
+    // Check results
+    for (uint32_t j = 0; j < n; ++j) {
+//      std::cout <<((uint32_t*)buf_data[1])[j] << "  ";
+      if(j >= i && j <= i+2) {
+       OCL_ASSERT(((uint32_t*)buf_data[1])[j] == color);
+      } else {
+       OCL_ASSERT(((uint32_t*)buf_data[1])[j] == 0);
+      }
+
+    }
+    memset(((uint32_t*)buf_data[1]), 0x0, sizeof(int)*n);
+    OCL_UNMAP_BUFFER(1);
+  }
+}
 MAKE_UTEST_FROM_FUNCTION(compiler_group_size1);
 MAKE_UTEST_FROM_FUNCTION(compiler_group_size2);
 MAKE_UTEST_FROM_FUNCTION(compiler_group_size3);
+MAKE_UTEST_FROM_FUNCTION(compiler_group_size4);