[Beignet] [PATCH 02/10] Utest: Add workgroup predicate tests

[grigore.lupescu at intel.com]

From: Grigore Lupescu <grigore.lupescu at intel.com>

Signed-off-by: Grigore Lupescu <grigore.lupescu at intel.com>
---
 kernels/compiler_workgroup_predicate.cl |  9 ++++++
 utests/CMakeLists.txt                   |  1 +
 utests/compiler_workgroup_predicate.cpp | 53 +++++++++++++++++++++++++++++++++
 3 files changed, 63 insertions(+)
 create mode 100644 kernels/compiler_workgroup_predicate.cl
 create mode 100644 utests/compiler_workgroup_predicate.cpp

diff --git a/kernels/compiler_workgroup_predicate.cl b/kernels/compiler_workgroup_predicate.cl
new file mode 100644
index 0000000..2b5f7c1
--- /dev/null
+++ b/kernels/compiler_workgroup_predicate.cl
@@ -0,0 +1,9 @@
+kernel void compiler_workgroup_all(global int *buf) {
+   uint gid = get_global_id(0);
+   buf[gid] = work_group_all(buf[gid]);
+}
+
+kernel void compiler_workgroup_any(global int *buf) {
+  uint gid = get_global_id(0);
+  buf[gid] = work_group_any(buf[gid]);
+}
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index c75e6b7..b849420 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -124,6 +124,7 @@ set (utests_sources
   compiler_atomic_functions.cpp
   compiler_async_copy.cpp
   compiler_workgroup_broadcast.cpp
+  compiler_workgroup_predicate.cpp
   compiler_workgroup_reduce.cpp
   compiler_async_stride_copy.cpp
   compiler_insn_selection_min.cpp
diff --git a/utests/compiler_workgroup_predicate.cpp b/utests/compiler_workgroup_predicate.cpp
new file mode 100644
index 0000000..d17e150
--- /dev/null
+++ b/utests/compiler_workgroup_predicate.cpp
@@ -0,0 +1,53 @@
+#include <cstdint>
+#include <cstring>
+#include <iostream>
+#include "utest_helper.hpp"
+
+typedef enum {TEST_WG_ALL, TEST_WG_ANY} TEST_WG_PREDICATE;
+
+static uint32_t test_array[60] = {
+    3, 5, 7, 9, 11, 3, 5, 7, 9, 11, 3, 5, 7, 9, 11, 3, 5, 7, 9, 11,
+    3, 5, 7, 9, 11, 3, 5, 7, 9, 11, 3, 5, 7, 9, 11, 3, 5, 7, 9, 11,
+    3, 5, 7, 9, 11, 3, 5, 7, 9, 11, 3, 5, 7, 9, 11, 3, 5, 7, 9, 11
+};
+
+void compiler_workgroup_predicate(TEST_WG_PREDICATE test_predicate)
+{
+  const size_t global_size = 60;
+  const size_t local_size = 60;
+
+  OCL_CREATE_BUFFER(buf[0], 0, global_size * sizeof(int32_t), NULL);
+  OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+  globals[0] = global_size;
+  locals[0] = local_size;
+
+  OCL_MAP_BUFFER(0);
+  memcpy(buf_data[0], &test_array[0], global_size * sizeof(int32_t));
+  OCL_UNMAP_BUFFER(0);
+
+  /* Run the kernel on GPU */
+  OCL_NDRANGE(1);
+
+  OCL_MAP_BUFFER(0);
+  for (int32_t i = 0; i < (int32_t) global_size; ++i) {
+    //printf("%d ", ((int32_t *)buf_data[0])[i]);
+    if(test_predicate == TEST_WG_ANY)
+        OCL_ASSERT(((uint32_t *)buf_data[0])[i] == 15);
+    if(test_predicate == TEST_WG_ALL)
+        OCL_ASSERT(((uint32_t *)buf_data[0])[i] == 1);
+  }
+  OCL_UNMAP_BUFFER(0);
+}
+
+void compiler_workgroup_all(void){
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_predicate", "compiler_workgroup_all");
+  compiler_workgroup_predicate(TEST_WG_ALL);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_all);
+
+void compiler_workgroup_any(void){
+  OCL_CREATE_KERNEL_FROM_FILE("compiler_workgroup_predicate", "compiler_workgroup_any");
+  compiler_workgroup_predicate(TEST_WG_ANY);
+}
+MAKE_UTEST_FROM_FUNCTION(compiler_workgroup_any);
-- 
2.5.0