diff options
| author | Pan Xiuli <xiuli.pan@intel.com> | 2016-10-19 14:37:19 +0800 |
|---|---|---|
| committer | Yang Rong <rong.r.yang@intel.com> | 2016-11-03 12:24:01 +0800 |
| commit | 2fc7c5249e09242802cb11b2d68eed9f59fb4b5c (patch) | |
| tree | 20e810a1951e1ba46a321f66ec40e85146153489 /kernels/compiler_subgroup_scan_exclusive.cl | |
| parent | f31ffa8e284c6fdd2a3ea478e932b7a6d80928ce (diff) | |
| download | beignet-2fc7c5249e09242802cb11b2d68eed9f59fb4b5c.tar.gz | |
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>
Diffstat (limited to 'kernels/compiler_subgroup_scan_exclusive.cl')
| -rw-r--r-- | kernels/compiler_subgroup_scan_exclusive.cl | 36 |
1 files changed, 36 insertions, 0 deletions
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); |