Utests: Add test cases for workgroup reduce max/min.

Signed-off-by: Junyan He <junyan.he@linux.intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

1 files changed, 28 insertions, 0 deletions

diff --git a/kernels/compiler_workgroup_reduce.cl b/kernels/compiler_workgroup_reduce.cl
new file mode 100644
index 00000000..27d306bc
--- /dev/null
+++ b/kernels/compiler_workgroup_reduce.cl
@@ -0,0 +1,28 @@
+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;
+}
+
+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;
+}
+
+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_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_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;
+}