Utest: Add workgroup scan exclusive tests

Added the following unit tests:
compiler_workgroup_scan_exclusive_add_int
compiler_workgroup_scan_exclusive_add_uint
compiler_workgroup_scan_exclusive_add_long
compiler_workgroup_scan_exclusive_add_ulong
compiler_workgroup_scan_exclusive_add_float
compiler_workgroup_scan_exclusive_max_int
compiler_workgroup_scan_exclusive_max_uint
compiler_workgroup_scan_exclusive_max_long
compiler_workgroup_scan_exclusive_max_ulong
compiler_workgroup_scan_exclusive_max_float
compiler_workgroup_scan_exclusive_min_int
compiler_workgroup_scan_exclusive_min_uint
compiler_workgroup_scan_exclusive_min_long
compiler_workgroup_scan_exclusive_min_ulong
compiler_workgroup_scan_exclusive_min_float

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

1 files changed, 98 insertions, 0 deletions

diff --git a/kernels/compiler_workgroup_scan_exclusive.cl b/kernels/compiler_workgroup_scan_exclusive.cl
new file mode 100644
index 00000000..14c1c618
--- /dev/null
+++ b/kernels/compiler_workgroup_scan_exclusive.cl
@@ -0,0 +1,98 @@
+/*
+ * Workgroup scan exclusive add functions
+ */
+kernel void compiler_workgroup_scan_exclusive_add_int(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int sum = work_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_uint(global uint *src, global uint *dst) {
+  uint val = src[get_global_id(0)];
+  uint sum = work_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_long(global long *src, global long *dst) {
+  long val = src[get_global_id(0)];
+  long sum = work_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_ulong(global ulong *src, global ulong *dst) {
+  ulong val = src[get_global_id(0)];
+  ulong sum = work_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_add_float(global float *src, global float *dst) {
+  float val = src[get_global_id(0)];
+  float sum = work_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+/*
+ * Workgroup scan exclusive max functions
+ */
+kernel void compiler_workgroup_scan_exclusive_max_int(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int sum = work_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_uint(global uint *src, global uint *dst) {
+  uint val = src[get_global_id(0)];
+  uint sum = work_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_long(global long *src, global long *dst) {
+  long val = src[get_global_id(0)];
+  long sum = work_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_ulong(global ulong *src, global ulong *dst) {
+  ulong val = src[get_global_id(0)];
+  ulong sum = work_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_max_float(global float *src, global float *dst) {
+  float val = src[get_global_id(0)];
+  float sum = work_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+/*
+ * Workgroup scan exclusive min functions
+ */
+kernel void compiler_workgroup_scan_exclusive_min_int(global int *src, global int *dst) {
+  int val = src[get_global_id(0)];
+  int sum = work_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_uint(global uint *src, global uint *dst) {
+  uint val = src[get_global_id(0)];
+  uint sum = work_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_long(global long *src, global long *dst) {
+  long val = src[get_global_id(0)];
+  long sum = work_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_ulong(global ulong *src, global ulong *dst) {
+  ulong val = src[get_global_id(0)];
+  ulong sum = work_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_workgroup_scan_exclusive_min_float(global float *src, global float *dst) {
+  float val = src[get_global_id(0)];
+  float sum = work_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}