Utest: Add test case for sub group short builtin functions

Signed-off-by: Pan Xiuli <xiuli.pan@intel.com>
Reviewed-by: Yang Rong <rong.r.yang@intel.com>

3 files changed, 94 insertions, 0 deletions

diff --git a/kernels/compiler_subgroup_reduce.cl b/kernels/compiler_subgroup_reduce.cl
index 6d7ecfd9..79d8e7d9 100644
--- a/kernels/compiler_subgroup_reduce.cl
+++ b/kernels/compiler_subgroup_reduce.cl
@@ -73,6 +73,17 @@ kernel void compiler_subgroup_reduce_add_float(global float *src, global float *
 /*
  * Subgroup reduce max functions
  */
+kernel void compiler_subgroup_reduce_max_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
+}
+kernel void compiler_subgroup_reduce_max_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  //printf("src is %d\n",val);
+  ushort sum = sub_group_reduce_max(val);
+  dst[get_global_id(0)] = sum;
+}
 kernel void compiler_subgroup_reduce_max_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_reduce_max(val);
@@ -106,6 +117,17 @@ kernel void compiler_subgroup_reduce_max_float(global float *src, global float *
 /*
  * Subgroup reduce min functions
  */
+kernel void compiler_subgroup_reduce_min_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
+kernel void compiler_subgroup_reduce_min_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  //printf("src is %d\n",val);
+  ushort sum = sub_group_reduce_min(val);
+  dst[get_global_id(0)] = sum;
+}
 kernel void compiler_subgroup_reduce_min_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_reduce_min(val);
diff --git a/kernels/compiler_subgroup_scan_exclusive.cl b/kernels/compiler_subgroup_scan_exclusive.cl
index ca0ada21..2c4b9285 100644
--- a/kernels/compiler_subgroup_scan_exclusive.cl
+++ b/kernels/compiler_subgroup_scan_exclusive.cl
@@ -2,6 +2,18 @@
  * Subgroup scan exclusive add functions
  */
 #ifndef HALF
+kernel void compiler_subgroup_scan_exclusive_add_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_exclusive_add_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = sub_group_scan_exclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
 kernel void compiler_subgroup_scan_exclusive_add_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_scan_exclusive_add(val);
@@ -35,6 +47,18 @@ kernel void compiler_subgroup_scan_exclusive_add_float(global float *src, global
 /*
  * Subgroup scan exclusive max functions
  */
+kernel void compiler_subgroup_scan_exclusive_max_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_exclusive_max_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = sub_group_scan_exclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
 kernel void compiler_subgroup_scan_exclusive_max_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_scan_exclusive_max(val);
@@ -68,6 +92,18 @@ kernel void compiler_subgroup_scan_exclusive_max_float(global float *src, global
 /*
  * Subgroup scan exclusive min functions
  */
+kernel void compiler_subgroup_scan_exclusive_min_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_exclusive_min_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = sub_group_scan_exclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
 kernel void compiler_subgroup_scan_exclusive_min_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_scan_exclusive_min(val);
diff --git a/kernels/compiler_subgroup_scan_inclusive.cl b/kernels/compiler_subgroup_scan_inclusive.cl
index e97521cf..def941c6 100644
--- a/kernels/compiler_subgroup_scan_inclusive.cl
+++ b/kernels/compiler_subgroup_scan_inclusive.cl
@@ -2,6 +2,18 @@
  * Subgroup scan inclusive add functions
  */
 #ifndef HALF
+kernel void compiler_subgroup_scan_inclusive_add_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_scan_inclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_inclusive_add_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = sub_group_scan_inclusive_add(val);
+  dst[get_global_id(0)] = sum;
+}
+
 kernel void compiler_subgroup_scan_inclusive_add_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_scan_inclusive_add(val);
@@ -35,6 +47,18 @@ kernel void compiler_subgroup_scan_inclusive_add_float(global float *src, global
 /*
  * Subgroup scan inclusive max functions
  */
+kernel void compiler_subgroup_scan_inclusive_max_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_scan_inclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_inclusive_max_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = sub_group_scan_inclusive_max(val);
+  dst[get_global_id(0)] = sum;
+}
+
 kernel void compiler_subgroup_scan_inclusive_max_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_scan_inclusive_max(val);
@@ -68,6 +92,18 @@ kernel void compiler_subgroup_scan_inclusive_max_float(global float *src, global
 /*
  * Subgroup scan inclusive min functions
  */
+kernel void compiler_subgroup_scan_inclusive_min_short(global short *src, global short *dst) {
+  short val = src[get_global_id(0)];
+  short sum = sub_group_scan_inclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
+kernel void compiler_subgroup_scan_inclusive_min_ushort(global ushort *src, global ushort *dst) {
+  ushort val = src[get_global_id(0)];
+  ushort sum = sub_group_scan_inclusive_min(val);
+  dst[get_global_id(0)] = sum;
+}
+
 kernel void compiler_subgroup_scan_inclusive_min_int(global int *src, global int *dst) {
   int val = src[get_global_id(0)];
   int sum = sub_group_scan_inclusive_min(val);