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_broadcast_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; /* set to 1 for debug, output of input-expected data */ #define DEBUG_STDOUT 0 @@ -54,24 +55,43 @@ static void compute_expected(WG_FUNCTION wg_func, else if(wg_func == WG_REDUCE_ADD) { T wg_sum = input[0]; - for(uint32_t i = 1; i < SIMD_SIZE; i++) - wg_sum += input[i]; + if(IS_HALF) { + float wg_sum_tmp = 0.0f; + for(uint32_t i = 0; i < SIMD_SIZE; i++) { + wg_sum_tmp += as_float(__half_to_float(input[i])); + } + wg_sum = __float_to_half(as_uint(wg_sum_tmp)); + } + else { + for(uint32_t i = 1; i < SIMD_SIZE; i++) + wg_sum += input[i]; + } for(uint32_t i = 0; i < SIMD_SIZE; i++) expected[i] = wg_sum; } else if(wg_func == WG_REDUCE_MAX) { T wg_max = input[0]; - for(uint32_t i = 1; i < SIMD_SIZE; i++) - wg_max = max(input[i], wg_max); + for(uint32_t i = 1; i < SIMD_SIZE; i++) { + if (IS_HALF) { + wg_max = (as_float(__half_to_float(input[i])) > as_float(__half_to_float(wg_max))) ? input[i] : wg_max; + } + else + wg_max = max(input[i], wg_max); + } for(uint32_t i = 0; i < SIMD_SIZE; i++) expected[i] = wg_max; } else if(wg_func == WG_REDUCE_MIN) { T wg_min = input[0]; - for(uint32_t i = 1; i < SIMD_SIZE; i++) - wg_min = min(input[i], wg_min); + for(uint32_t i = 1; i < SIMD_SIZE; i++) { + if (IS_HALF) { + wg_min= (as_float(__half_to_float(input[i])) < as_float(__half_to_float(wg_min))) ? input[i] : wg_min; + } + else + wg_min = min(input[i], wg_min); + } for(uint32_t i = 0; i < SIMD_SIZE; i++) expected[i] = wg_min; } @@ -115,6 +135,8 @@ static void generate_data(WG_FUNCTION wg_func, /* add trailing random bits, tests GENERAL cases */ input[gid + lid] += (rand() % 112); /* always last bit is 1, ideal test ALL/ANY */ + if (IS_HALF) + input[gid + lid] = __float_to_half(as_uint((float)input[gid + lid]/2)); } else { input[gid + lid] += rand(); input[gid + lid] += rand() / ((float)RAND_MAX + 1); @@ -185,8 +207,22 @@ static void subgroup_generic(WG_FUNCTION wg_func, for (uint32_t i = 0; i < WG_GLOBAL_SIZE; i++) if(((T *)buf_data[1])[i] != *(expected + i)) { + if (IS_HALF) { + float num_computed = as_float(__half_to_float(((T *)buf_data[1])[i])); + float num_expected = as_float(__half_to_float(*(expected + i))); + float num_diff = abs(num_computed - num_expected) / abs(num_expected); + if (num_diff > 0.03f) { + mismatches++; + } +#if DEBUG_STDOUT + /* output mismatch */ + cout << "Err at " << i << ", " << num_computed + << " != " << num_expected << " diff: " <<num_diff <<endl; +#endif + //} + } /* found mismatch on integer, increment */ - if (numeric_limits<T>::is_integer) { + else if (numeric_limits<T>::is_integer) { mismatches++; #if DEBUG_STDOUT @@ -305,6 +341,22 @@ void compiler_subgroup_reduce_add_float(void) subgroup_generic(WG_REDUCE_ADD, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_add_float); +void compiler_subgroup_reduce_add_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_reduce.cl", + "compiler_subgroup_reduce_add_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_REDUCE_ADD, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_add_half); /* * Workgroup reduce max utest functions @@ -364,6 +416,22 @@ void compiler_subgroup_reduce_max_float(void) subgroup_generic(WG_REDUCE_MAX, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_max_float); +void compiler_subgroup_reduce_max_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_reduce.cl", + "compiler_subgroup_reduce_max_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_REDUCE_MAX, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_max_half); /* * Workgroup reduce min utest functions @@ -423,3 +491,19 @@ void compiler_subgroup_reduce_min_float(void) subgroup_generic(WG_REDUCE_MIN, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_min_float); +void compiler_subgroup_reduce_min_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_reduce.cl", + "compiler_subgroup_reduce_min_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_REDUCE_MIN, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_reduce_min_half); diff --git a/utests/compiler_subgroup_scan_exclusive.cpp b/utests/compiler_subgroup_scan_exclusive.cpp index 1a21b59..dc3b15d 100644 --- a/utests/compiler_subgroup_scan_exclusive.cpp +++ b/utests/compiler_subgroup_scan_exclusive.cpp @@ -9,6 +9,7 @@ #include "utest_helper.hpp" using namespace std; +static bool IS_HALF = false; /* set to 1 for debug, output of input-expected data */ #define DEBUG_STDOUT 0 @@ -38,30 +39,49 @@ static void compute_expected(WG_FUNCTION wg_func, { expected[0] = 0; expected[1] = input[0]; - for(uint32_t i = 2; i < SIMD_SIZE; i++) - expected[i] = input[i - 1] + expected[i - 1]; + for(uint32_t i = 2; i < SIMD_SIZE; i++) { + if (IS_HALF) + expected[i] = __float_to_half(as_uint(as_float(__half_to_float(input[i - 1])) + + as_float(__half_to_float(expected[i - 1])))); + else + expected[i] = input[i - 1] + expected[i - 1]; + } } else if(wg_func == WG_SCAN_EXCLUSIVE_MAX) { - if(numeric_limits<T>::is_integer) + if(IS_HALF) + expected[0] = 0xFC00; + else if(numeric_limits<T>::is_integer) expected[0] = numeric_limits<T>::min(); else expected[0] = - numeric_limits<T>::infinity(); expected[1] = input[0]; - for(uint32_t i = 2; i < SIMD_SIZE; i++) - expected[i] = max(input[i - 1], expected[i - 1]); + for(uint32_t i = 2; i < SIMD_SIZE; i++) { + if (IS_HALF) + expected[i] = (as_float(__half_to_float(input[i - 1])) > as_float(__half_to_float(expected[i - 1]))) ? + input[i - 1] : expected[i - 1]; + else + expected[i] = max(input[i - 1], expected[i - 1]); + } } else if(wg_func == WG_SCAN_EXCLUSIVE_MIN) { - if(numeric_limits<T>::is_integer) + if(IS_HALF) + expected[0] = 0x7C00; + else if(numeric_limits<T>::is_integer) expected[0] = numeric_limits<T>::max(); else expected[0] = numeric_limits<T>::infinity(); expected[1] = input[0]; - for(uint32_t i = 2; i < SIMD_SIZE; i++) - expected[i] = min(input[i - 1], expected[i - 1]); + for(uint32_t i = 2; i < SIMD_SIZE; i++) { + if (IS_HALF) + expected[i] = (as_float(__half_to_float(input[i - 1])) < as_float(__half_to_float(expected[i - 1]))) ? + input[i - 1] : expected[i - 1]; + else + expected[i] = min(input[i - 1], expected[i - 1]); + } } } @@ -101,6 +121,8 @@ static void generate_data(WG_FUNCTION wg_func, input[gid + lid] += ((rand() % 2 - 1) * base_val); /* add trailing random bits, tests GENERAL cases */ input[gid + lid] += (rand() % 112); + if (IS_HALF) + input[gid + lid] = __float_to_half(as_uint((float)input[gid + lid]/2)); #if DEBUG_STDOUT /* output generated input */ @@ -166,8 +188,21 @@ static void subgroup_generic(WG_FUNCTION wg_func, for (uint32_t i = 0; i < WG_GLOBAL_SIZE; i++) if(((T *)buf_data[1])[i] != *(expected + i)) { + if (IS_HALF) { + float num_computed = as_float(__half_to_float(((T *)buf_data[1])[i])); + float num_expected = as_float(__half_to_float(*(expected + i))); + float num_diff = abs(num_computed - num_expected) / abs(num_expected); + if (num_diff > 0.03f) { + mismatches++; +#if DEBUG_STDOUT + /* output mismatch */ + cout << "Err at " << i << ", " << num_computed + << " != " << num_expected <<" diff: " <<num_diff <<endl; +#endif + } + } /* found mismatch on integer, increment */ - if(numeric_limits<T>::is_integer){ + else if (numeric_limits<T>::is_integer) { mismatches++; #if DEBUG_STDOUT @@ -261,6 +296,22 @@ void compiler_subgroup_scan_exclusive_add_float(void) subgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_add_float); +void compiler_subgroup_scan_exclusive_add_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_scan_exclusive.cl", + "compiler_subgroup_scan_exclusive_add_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_SCAN_EXCLUSIVE_ADD, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_add_half); /* * Workgroup scan_exclusive max utest functions @@ -320,6 +371,22 @@ void compiler_subgroup_scan_exclusive_max_float(void) subgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_max_float); +void compiler_subgroup_scan_exclusive_max_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_scan_exclusive.cl", + "compiler_subgroup_scan_exclusive_max_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_SCAN_EXCLUSIVE_MAX, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_max_half); /* * Workgroup scan_exclusive min utest functions @@ -379,3 +446,19 @@ void compiler_subgroup_scan_exclusive_min_float(void) subgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_min_float); +void compiler_subgroup_scan_exclusive_min_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_scan_exclusive.cl", + "compiler_subgroup_scan_exclusive_min_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_SCAN_EXCLUSIVE_MIN, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_exclusive_min_half); diff --git a/utests/compiler_subgroup_scan_inclusive.cpp b/utests/compiler_subgroup_scan_inclusive.cpp index fa32855..1c8fd8a 100644 --- a/utests/compiler_subgroup_scan_inclusive.cpp +++ b/utests/compiler_subgroup_scan_inclusive.cpp @@ -9,6 +9,7 @@ #include "utest_helper.hpp" using namespace std; +static bool IS_HALF = false; /* set to 1 for debug, output of input-expected data */ #define DEBUG_STDOUT 0 @@ -37,20 +38,35 @@ static void compute_expected(WG_FUNCTION wg_func, if(wg_func == WG_SCAN_INCLUSIVE_ADD) { expected[0] = input[0]; - for(uint32_t i = 1; i < SIMD_SIZE; i++) - expected[i] = input[i] + expected[i - 1]; + for(uint32_t i = 1; i < SIMD_SIZE; i++) { + if (IS_HALF) + expected[i] = __float_to_half(as_uint(as_float(__half_to_float(input[i])) + + as_float(__half_to_float(expected[i - 1])))); + else + expected[i] = input[i] + expected[i - 1]; + } } else if(wg_func == WG_SCAN_INCLUSIVE_MAX) { expected[0] = input[0]; - for(uint32_t i = 1; i < SIMD_SIZE; i++) - expected[i] = max(input[i], expected[i - 1]); + for(uint32_t i = 1; i < SIMD_SIZE; i++) { + if (IS_HALF) + expected[i] = (as_float(__half_to_float(input[i])) > as_float(__half_to_float(expected[i - 1]))) ? + input[i] : expected[i - 1]; + else + expected[i] = max(input[i], expected[i - 1]); + } } else if(wg_func == WG_SCAN_INCLUSIVE_MIN) { expected[0] = input[0]; - for(uint32_t i = 1; i < SIMD_SIZE; i++) - expected[i] = min(input[i], expected[i - 1]); + for(uint32_t i = 1; i < SIMD_SIZE; i++) { + if (IS_HALF) + expected[i] = (as_float(__half_to_float(input[i])) < as_float(__half_to_float(expected[i - 1]))) ? + input[i] : expected[i - 1]; + else + expected[i] = min(input[i], expected[i - 1]); + } } } @@ -91,6 +107,8 @@ static void generate_data(WG_FUNCTION wg_func, input[gid + lid] += ((rand() % 2 - 1) * base_val); /* add trailing random bits, tests GENERAL cases */ input[gid + lid] += (rand() % 112); + if (IS_HALF) + input[gid + lid] = __float_to_half(as_uint((float)input[gid + lid]/2)); #if DEBUG_STDOUT /* output generated input */ @@ -156,8 +174,21 @@ static void subgroup_generic(WG_FUNCTION wg_func, for (uint32_t i = 0; i < WG_GLOBAL_SIZE; i++) if(((T *)buf_data[1])[i] != *(expected + i)) { + if (IS_HALF) { + float num_computed = as_float(__half_to_float(((T *)buf_data[1])[i])); + float num_expected = as_float(__half_to_float(*(expected + i))); + float num_diff = abs(num_computed - num_expected) / abs(num_expected); + if (num_diff > 0.03f) { + mismatches++; +#if DEBUG_STDOUT + /* output mismatch */ + cout << "Err at " << i << ", " << num_computed + << " != " << num_expected <<" diff: " <<num_diff <<endl; +#endif + } + } /* found mismatch on integer, increment */ - if(numeric_limits<T>::is_integer){ + else if (numeric_limits<T>::is_integer) { mismatches++; #if DEBUG_STDOUT @@ -251,6 +282,22 @@ void compiler_subgroup_scan_inclusive_add_float(void) subgroup_generic(WG_SCAN_INCLUSIVE_ADD, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_add_float); +void compiler_subgroup_scan_inclusive_add_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_scan_inclusive.cl", + "compiler_subgroup_scan_inclusive_add_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_SCAN_INCLUSIVE_ADD, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_add_half); /* * Workgroup scan_inclusive max utest functions @@ -310,6 +357,22 @@ void compiler_subgroup_scan_inclusive_max_float(void) subgroup_generic(WG_SCAN_INCLUSIVE_MAX, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_max_float); +void compiler_subgroup_scan_inclusive_max_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_scan_inclusive.cl", + "compiler_subgroup_scan_inclusive_max_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_SCAN_INCLUSIVE_MAX, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_max_half); /* * Workgroup scan_inclusive min utest functions @@ -369,4 +432,19 @@ void compiler_subgroup_scan_inclusive_min_float(void) subgroup_generic(WG_SCAN_INCLUSIVE_MIN, input, expected); } MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_min_float); - +void compiler_subgroup_scan_inclusive_min_half(void) +{ + if(!cl_check_subgroups()) + return; + if(!cl_check_half()) + return; + cl_half *input = NULL; + cl_half *expected = NULL; + IS_HALF = true; + OCL_CALL(cl_kernel_init, "compiler_subgroup_scan_inclusive.cl", + "compiler_subgroup_scan_inclusive_min_half", + SOURCE, "-DHALF"); + subgroup_generic(WG_SCAN_INCLUSIVE_MIN, input, expected); + IS_HALF = false; +} +MAKE_UTEST_FROM_FUNCTION(compiler_subgroup_scan_inclusive_min_half); -- 2.5.0 _______________________________________________ Beignet mailing list [email protected] https://lists.freedesktop.org/mailman/listinfo/beignet
