Utest: Add workgroup reduce any/all tests

Added the following unit tests:
compiler_workgroup_any
compiler_workgroup_all
compiler_workgroup_reduce_add_int
compiler_workgroup_reduce_add_uint
compiler_workgroup_reduce_add_long
compiler_workgroup_reduce_add_ulong
compiler_workgroup_reduce_add_float
compiler_workgroup_reduce_max_int
compiler_workgroup_reduce_max_uint
compiler_workgroup_reduce_max_long
compiler_workgroup_reduce_max_ulong
compiler_workgroup_reduce_max_float
compiler_workgroup_reduce_min_int
compiler_workgroup_reduce_min_uint
compiler_workgroup_reduce_min_long
compiler_workgroup_reduce_min_ulong
compiler_workgroup_reduce_min_float

Signed-off-by: Grigore Lupescu <grigore.lupescu at intel.com>
Reviewed-by: Pan Xiuli <xiuli.pan@intel.com>

1 files changed, 122 insertions, 25 deletions

diff --git a/kernels/compiler_workgroup_reduce.cl b/kernels/compiler_workgroup_reduce.cl
index 1fc57b51..d2f37ca9 100644
--- a/kernels/compiler_workgroup_reduce.cl
+++ b/kernels/compiler_workgroup_reduce.cl
@@ -1,40 +1,137 @@
-kernel void compiler_workgroup_reduce_min_uniform(uint src, global uint *dst) {
-   uint min_val = work_group_reduce_min(src);
-   dst[get_local_id(0)] = min_val;
+/*
+ * Workgroup any all functions
+ */
+kernel void compiler_workgroup_any(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int predicate = work_group_any(val);
+  dst[get_global_id(0)] = predicate;
+}
+kernel void compiler_workgroup_all(global int *src, global int *dst) {
+  char val = src[get_global_id(0)];
+  int predicate = work_group_all(val);
+  dst[get_global_id(0)] = predicate;
 }
 
-kernel void compiler_workgroup_reduce_min_uint(global uint *src, global uint *dst) {
-   uint val = src[get_local_id(0)];
-   uint min_val = work_group_reduce_min(val);
-   dst[get_local_id(0)] = min_val;
+/*
+ * Workgroup reduce add functions
+ */
+kernel void compiler_workgroup_reduce_add_char(global char *src, global char *dst) {
+  char val = src[get_global_id(0)];
+  char sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
 }
 
-kernel void compiler_workgroup_reduce_max_uint(global uint *src, global uint *dst) {
-   uint val = src[get_local_id(0)];
-   uint max_val = work_group_reduce_max(val);
-   dst[get_local_id(0)] = max_val;
+kernel void compiler_workgroup_reduce_add_uchar(global uchar *src, global uchar *dst) {
+  uchar val = src[get_global_id(0)];
+  uchar sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
 }
 
-kernel void compiler_workgroup_reduce_min_float(global float *src, global float *dst) {
-   float val = src[get_local_id(0)];
-   float min_val = work_group_reduce_min(val);
-   dst[get_local_id(0)] = min_val;
+kernel void compiler_workgroup_reduce_add_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
 }
 
-kernel void compiler_workgroup_reduce_max_float(global float *src, global float *dst) {
-   float val = src[get_local_id(0)];
-   float max_val = work_group_reduce_max(val);
-   dst[get_local_id(0)] = max_val;
+kernel void compiler_workgroup_reduce_add_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_add_int(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
 }
 
 kernel void compiler_workgroup_reduce_add_uint(global uint *src, global uint *dst) {
-   uint val = src[get_local_id(0)];
-   uint sum = work_group_reduce_add(val);
-   dst[get_local_id(0)] = sum;
+  uint val = src[get_global_id(0)];
+  uint sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_add_long(global long *src, global long *dst) {
+  long val = src[get_global_id(0)];
+  long sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_add_ulong(global ulong *src, global ulong *dst) {
+  ulong val = src[get_global_id(0)];
+  ulong sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
 }
 
 kernel void compiler_workgroup_reduce_add_float(global float *src, global float *dst) {
-   float val = src[get_local_id(0)];
-   float sum = work_group_reduce_add(val);
-   dst[get_local_id(0)] = sum;
+  float val = src[get_global_id(0)];
+  float sum = work_group_reduce_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+/*
+ * Workgroup reduce max functions
+ */
+kernel void compiler_workgroup_reduce_max_int(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int sum = work_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_max_uint(global uint *src, global uint *dst) {
+  uint val = src[get_global_id(0)];
+  uint sum = work_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_max_long(global long *src, global long *dst) {
+  long val = src[get_global_id(0)];
+  long sum = work_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
 }
+
+kernel void compiler_workgroup_reduce_max_ulong(global ulong *src, global ulong *dst) {
+  ulong val = src[get_global_id(0)];
+  ulong sum = work_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_max_float(global float *src, global float *dst) {
+  float val = src[get_global_id(0)];
+  float sum = work_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+/*
+ * Workgroup reduce min functions
+ */
+kernel void compiler_workgroup_reduce_min_int(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int sum = work_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_min_uint(global uint *src, global uint *dst) {
+  uint val = src[get_global_id(0)];
+  uint sum = work_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_min_long(global long *src, global long *dst) {
+  long val = src[get_global_id(0)];
+  long sum = work_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_min_ulong(global ulong *src, global ulong *dst) {
+  ulong val = src[get_global_id(0)];
+  ulong sum = work_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_reduce_min_float(global float *src, global float *dst) {
+  float val = src[get_global_id(0)];
+  float sum = work_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
+