[Beignet] [PATCH V2 7/9] Utest: Add test for half type subgroup functions
Yang, Rong R
rong.r.yang at intel.com
Fri Aug 12 07:30:53 UTC 2016
One comment.
> -----Original Message-----
> From: Beignet [mailto:beignet-bounces at lists.freedesktop.org] On Behalf Of
> Xiuli Pan
> Sent: Monday, August 8, 2016 11:31
> To: beignet at lists.freedesktop.org
> Cc: Pan, Xiuli <xiuli.pan at intel.com>
> Subject: [Beignet] [PATCH V2 7/9] Utest: Add test for half type subgroup
> functions
>
> From: Pan Xiuli <xiuli.pan at intel.com>
>
> Check if device support subgroup and half first, use build options to hide
> code for unsported device.
>
> Signed-off-by: Pan Xiuli <xiuli.pan at intel.com>
> ---
> 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
> Beignet at lists.freedesktop.org
> https://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list