diff options
| author | Pan Xiuli <xiuli.pan@intel.com> | 2016-08-18 12:56:37 +0800 |
|---|---|---|
| committer | Yang Rong <rong.r.yang@intel.com> | 2016-08-31 16:23:42 +0800 |
| commit | 150b39e35c3f5b3bb07ab7a542eeed7ebeff57df (patch) | |
| tree | bb9ccc50db5ce5ea1fa3e27d46da8e98839f012a /kernels/compiler_subgroup_scan_exclusive.cl | |
| parent | 7d00b274eabf4f31be61e6a7909ed83d7ffbd142 (diff) | |
| download | beignet-150b39e35c3f5b3bb07ab7a542eeed7ebeff57df.tar.gz | |
Utest: Add test for half type subgroup functions
Check if device support subgroup and half first, use build options
to hide code for unsported device.
V2: Fix half part test case for utest multithread.
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 | 19 |
1 files changed, 19 insertions, 0 deletions
diff --git a/kernels/compiler_subgroup_scan_exclusive.cl b/kernels/compiler_subgroup_scan_exclusive.cl index afc00d0e..ca0ada21 100644 --- a/kernels/compiler_subgroup_scan_exclusive.cl +++ b/kernels/compiler_subgroup_scan_exclusive.cl @@ -1,6 +1,7 @@ /* * Subgroup scan exclusive add functions */ +#ifndef HALF 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); @@ -96,3 +97,21 @@ kernel void compiler_subgroup_scan_exclusive_min_float(global float *src, global float sum = sub_group_scan_exclusive_min(val); dst[get_global_id(0)] = sum; } +#else +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +kernel void compiler_subgroup_scan_exclusive_add_half(global half *src, global half *dst) { + half val = src[get_global_id(0)]; + half sum = sub_group_scan_exclusive_add(val); + dst[get_global_id(0)] = sum; +} +kernel void compiler_subgroup_scan_exclusive_max_half(global half *src, global half *dst) { + half val = src[get_global_id(0)]; + half sum = sub_group_scan_exclusive_max(val); + dst[get_global_id(0)] = sum; +} +kernel void compiler_subgroup_scan_exclusive_min_half(global half *src, global half *dst) { + half val = src[get_global_id(0)]; + half sum = sub_group_scan_exclusive_min(val); + dst[get_global_id(0)] = sum; +} +#endif |