One comment. > -----Original Message----- > From: Beignet [mailto:[email protected]] On Behalf Of > Xiuli Pan > Sent: Monday, August 8, 2016 11:31 > To: [email protected] > Cc: Pan, Xiuli <[email protected]> > Subject: [Beignet] [PATCH V2 7/9] Utest: Add test for half type subgroup > functions > > From: Pan Xiuli <[email protected]> > > Check if device support subgroup and half first, use build options to hide > code for unsported device. > > Signed-off-by: Pan Xiuli <[email protected]> > --- > kernels/compiler_subgroup_broadcast.cl | 16 ++++- > kernels/compiler_subgroup_reduce.cl | 19 ++++++ > kernels/compiler_subgroup_scan_exclusive.cl | 19 ++++++ > kernels/compiler_subgroup_scan_inclusive.cl | 19 ++++++ > utests/compiler_subgroup_broadcast.cpp | 27 ++++++-- > utests/compiler_subgroup_reduce.cpp | 98 > +++++++++++++++++++++++++-- > utests/compiler_subgroup_scan_exclusive.cpp | 101 > +++++++++++++++++++++++++--- > utests/compiler_subgroup_scan_inclusive.cpp | 94 > +++++++++++++++++++++++--- > 8 files changed, 364 insertions(+), 29 deletions(-) > > diff --git a/kernels/compiler_subgroup_broadcast.cl > b/kernels/compiler_subgroup_broadcast.cl > index 4f21cf5..8c155ee 100644 > --- a/kernels/compiler_subgroup_broadcast.cl > +++ b/kernels/compiler_subgroup_broadcast.cl > @@ -1,7 +1,7 @@ > /* > * Subgroup broadcast 1D functions > */ > - > +#ifndef HALF > kernel void compiler_subgroup_broadcast_imm_int(global int *src, > global int *dst, > uint simd_id) @@ -32,3 > +32,17 @@ kernel void > compiler_subgroup_broadcast_long(global long *src, > long broadcast_val = sub_group_broadcast(val, simd_id); > dst[index] = broadcast_val; > } > +#else > +#pragma OPENCL EXTENSION cl_khr_fp16 : enable kernel void > +compiler_subgroup_broadcast_half(global half *src, > + global half *dst, > + uint simd_id) { > + uint index = get_global_id(0); > + > + half val = src[index]; > + half broadcast_val = sub_group_broadcast(val, simd_id); > + printf("%d val %d is %d\n",index,as_ushort(val), > +as_ushort(broadcast_val)); > + dst[index] = broadcast_val; > +} > +#endif > diff --git a/kernels/compiler_subgroup_reduce.cl > b/kernels/compiler_subgroup_reduce.cl > index 77ffb07..6d7ecfd 100644 > --- a/kernels/compiler_subgroup_reduce.cl > +++ b/kernels/compiler_subgroup_reduce.cl > @@ -1,6 +1,7 @@ > /* > * Subgroup any all functions > */ > +#ifndef HALF > kernel void compiler_subgroup_any(global int *src, global int *dst) { > int val = src[get_global_id(0)]; > int predicate = sub_group_any(val); > @@ -134,3 +135,21 @@ kernel void > compiler_subgroup_reduce_min_float(global float *src, global float * > float sum = sub_group_reduce_min(val); > dst[get_global_id(0)] = sum; > } > +#else > +#pragma OPENCL EXTENSION cl_khr_fp16 : enable kernel void > +compiler_subgroup_reduce_add_half(global half *src, global half *dst) { > + half val = src[get_global_id(0)]; > + half sum = sub_group_reduce_add(val); > + dst[get_global_id(0)] = sum; > +} > +kernel void compiler_subgroup_reduce_max_half(global half *src, global > +half *dst) { > + half val = src[get_global_id(0)]; > + half sum = sub_group_reduce_max(val); > + dst[get_global_id(0)] = sum; > +} > +kernel void compiler_subgroup_reduce_min_half(global half *src, global > +half *dst) { > + half val = src[get_global_id(0)]; > + half sum = sub_group_reduce_min(val); > + dst[get_global_id(0)] = sum; > +} > +#endif > diff --git a/kernels/compiler_subgroup_scan_exclusive.cl > b/kernels/compiler_subgroup_scan_exclusive.cl > index afc00d0..ca0ada2 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 > diff --git a/kernels/compiler_subgroup_scan_inclusive.cl > b/kernels/compiler_subgroup_scan_inclusive.cl > index da1a6e6..e97521c 100644 > --- a/kernels/compiler_subgroup_scan_inclusive.cl > +++ b/kernels/compiler_subgroup_scan_inclusive.cl > @@ -1,6 +1,7 @@ > /* > * Subgroup scan inclusive add functions > */ > +#ifndef HALF > 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); > @@ -96,3 +97,21 @@ kernel void > compiler_subgroup_scan_inclusive_min_float(global float *src, global > float sum = sub_group_scan_inclusive_min(val); > dst[get_global_id(0)] = sum; > } > +#else > +#pragma OPENCL EXTENSION cl_khr_fp16 : enable kernel void > +compiler_subgroup_scan_inclusive_add_half(global half *src, global half > +*dst) { > + half val = src[get_global_id(0)]; > + half sum = sub_group_scan_inclusive_add(val); > + dst[get_global_id(0)] = sum; > +} > +kernel void compiler_subgroup_scan_inclusive_max_half(global half *src, > +global half *dst) { > + half val = src[get_global_id(0)]; > + half sum = sub_group_scan_inclusive_max(val); > + dst[get_global_id(0)] = sum; > +} > +kernel void compiler_subgroup_scan_inclusive_min_half(global half *src, > +global half *dst) { > + half val = src[get_global_id(0)]; > + half sum = sub_group_scan_inclusive_min(val); > + dst[get_global_id(0)] = sum; > +} > +#endif > diff --git a/utests/compiler_subgroup_broadcast.cpp > b/utests/compiler_subgroup_broadcast.cpp > index 2835161..9a7979c 100644 > --- a/utests/compiler_subgroup_broadcast.cpp > +++ b/utests/compiler_subgroup_broadcast.cpp > @@ -59,10 +59,15 @@ static void generate_data(T* &input, > /* initially 0, augment after */ > input[gid + lid] = 0; > > - /* check all data types, test ideal for QWORD types */ > - input[gid + lid] += ((rand() % 2 - 1) * base_val); > - /* add trailing random bits, tests GENERAL cases */ > - input[gid + lid] += (rand() % 112); > + if(sizeof(T) == 2) { > + input[gid + lid] = __float_to_half(as_uint((float)(gid + lid))); > + } > + else { > + /* check all data types, test ideal for QWORD types */ > + input[gid + lid] += ((rand() % 2 - 1) * base_val); > + /* add trailing random bits, tests GENERAL cases */ > + input[gid + lid] += (rand() % 112); > + } > > #if DEBUG_STDOUT > /* output generated input */ > @@ -185,3 +190,17 @@ void compiler_subgroup_broadcast_long(void) > subgroup_generic(input, expected); > } > > MAKE_UTEST_FROM_FUNCTION_WITH_ISSUE(compiler_subgroup_broadca > st_long); > +void compiler_subgroup_broadcast_half(void) > +{ > + if(!cl_check_subgroups()) > + return; > + if(!cl_check_half()) > + return; > + cl_half *input = NULL; > + cl_half *expected = NULL; > + OCL_CALL(cl_kernel_init, "compiler_subgroup_broadcast.cl", > + "compiler_subgroup_broadcast_half", > + SOURCE, "-DHALF"); > + subgroup_generic(input, expected); > +} > +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_broadcast_half); > diff --git a/utests/compiler_subgroup_reduce.cpp > b/utests/compiler_subgroup_reduce.cpp > index 3c3df06..9ab1d75 100644 > --- a/utests/compiler_subgroup_reduce.cpp > +++ b/utests/compiler_subgroup_reduce.cpp > @@ -9,6 +9,7 @@ > #include "utest_helper.hpp" > > using namespace std; > +static bool IS_HALF = false; static var without lock is not thread safe, right?
> _______________________________________________ > Beignet mailing list > [email protected] > https://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
