From aa2e9a7cb9d9cb8d92ca645a1d94f45a128d61ce Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Mon, 13 Mar 2023 09:44:23 -0400 Subject: [PATCH 1/8] [SYCL] Add marray support to common + some math functions This patch adds marray support to all functions from Table 179 of SYCL 2020 spec + to functions fabs, ilogb, fmax, fmin, ldexp, pown, rootn from Table 175 + to function exp10 from Table 177. --- sycl/include/sycl/builtins.hpp | 185 +++++++++++++++++++++++++++++++++ 1 file changed, 185 insertions(+) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 799970e256dcc..2b855f316f951 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -122,10 +122,27 @@ __SYCL_MATH_FUNCTION_OVERLOAD_FM(log2) __SYCL_MATH_FUNCTION_OVERLOAD_FM(log10) __SYCL_MATH_FUNCTION_OVERLOAD_FM(sqrt) __SYCL_MATH_FUNCTION_OVERLOAD_FM(rsqrt) +__SYCL_MATH_FUNCTION_OVERLOAD_FM(fabs) #undef __SYCL_MATH_FUNCTION_OVERLOAD_FM #undef __SYCL_MATH_FUNCTION_OVERLOAD_IMPL +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + ilogb(marray x) __NOEXC { + marray res; + for (size_t i = 0; i < N / 2; i++) { + vec partial_res = + __sycl_std::__invoke_ilogb>(detail::to_vec2(x, i * 2)); + std::memcpy(&res[i * 2], &partial_res, sizeof(vec)); + } + if (N % 2) { + res[N - 1] = __sycl_std::__invoke_ilogb(x[N - 1]); + } + return res; +} + #define __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(NAME) \ marray res; \ for (size_t i = 0; i < N / 2; i++) { \ @@ -170,6 +187,100 @@ inline __SYCL_ALWAYS_INLINE #undef __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL +#define __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(NAME) \ + template \ + inline __SYCL_ALWAYS_INLINE \ + std::enable_if_t::value, marray> \ + NAME(marray x, T y) __NOEXC { \ + marray res; \ + sycl::vec y_vec{y, y}; \ + for (size_t i = 0; i < N / 2; i++) { \ + auto partial_res = __sycl_std::__invoke_##NAME>( \ + detail::to_vec2(x, i * 2), y_vec); \ + std::memcpy(&res[i * 2], &partial_res, sizeof(vec)); \ + } \ + if (N % 2) { \ + res[N - 1] = __sycl_std::__invoke_##NAME(x[N - 1], y_vec[0]); \ + } \ + return res; \ + } + +__SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmax) + __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) + +#undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD + + template + inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> ldexp( + marray x, marray k) __NOEXC { + marray res; + for (size_t i = 0; i < N; i++) { + res[i] = __sycl_std::__invoke_ldexp(x[i], k[i]); + } + return res; +} + +#undef __SYCL_MATH_FUNCTION_2_GENINT_K_OVERLOAD_IMPL + +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + ldexp(marray x, int k) __NOEXC { + marray res; + for (size_t i = 0; i < N; i++) { + res[i] = __sycl_std::__invoke_ldexp(x[i], k); + } + return res; +} + +#undef __SYCL_MATH_FUNCTION_2_INT_K_OVERLOAD_IMPL + +#define __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(NAME) \ + marray res; \ + for (size_t i = 0; i < N; i++) { \ + res[i] = __sycl_std::__invoke_##NAME(x[i], y[i]); \ + } \ + return res; + +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + pown(marray x, marray y) __NOEXC { + __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(pown) +} + +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + rootn(marray x, marray y) __NOEXC { + __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(rootn) +} + +#undef __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL + +#define __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(NAME) \ + marray res; \ + for (size_t i = 0; i < N; i++) { \ + res[i] = __sycl_std::__invoke_##NAME(x[i], y); \ + } \ + return res; + +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + pown(marray x, int y) __NOEXC { + __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(pown) +} + +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + rootn(marray x, + int y) __NOEXC{__SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(rootn)} + +#undef __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL + #define __SYCL_MATH_FUNCTION_3_OVERLOAD(NAME) \ template \ inline __SYCL_ALWAYS_INLINE \ @@ -789,6 +900,79 @@ detail::enable_if_t::value, T> sign(T x) __NOEXC { return __sycl_std::__invoke_sign(x); } +// marray common functions + +// TODO: can be optimized in the way math functions are optimized (usage of +// vec) +#define __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(NAME, ARGTYPE, ARGVAL, \ + ITEM) \ + template ::value>> \ + sycl::marray, T::size()> NAME(ARGTYPE ARGVAL) \ + __NOEXC { \ + sycl::marray, T::size()> res; \ + for (int i = 0; i < ARGVAL.size(); i++) { \ + res[i] = NAME(ITEM); \ + } \ + return res; \ + } + +#define __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(NAME, ARG1, ARG2, ITEM1, \ + ITEM2) \ + template ::value>> \ + sycl::marray, T::size()> NAME(ARG1, ARG2) \ + __NOEXC { \ + sycl::marray, T::size()> res; \ + for (int i = 0; i < x.size(); i++) { \ + res[i] = NAME(ITEM1, ITEM2); \ + } \ + return res; \ + } + +#define __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(NAME, ARG1, ARG2, ARG3, \ + ITEM1, ITEM2, ITEM3) \ + template ::value>> \ + sycl::marray, T::size()> NAME(ARG1, ARG2, \ + ARG3) __NOEXC { \ + sycl::marray, T::size()> res; \ + for (int i = 0; i < x.size(); i++) { \ + res[i] = NAME(ITEM1, ITEM2, ITEM3); \ + } \ + return res; \ + } + +__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(clamp, T x, T minval, T maxval, + x[i], minval[i], maxval[i]) +__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD( + clamp, T x, detail::marray_element_type minval, + detail::marray_element_type maxval, x[i], minval, maxval) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(degrees, T, radians, radians[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, T y, x[i], y[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, + detail::marray_element_type y, + x[i], y) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(min, T x, T y, x[i], y[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(min, T x, + detail::marray_element_type y, + x[i], y) +__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, T a, x[i], y[i], + a[i]) +__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, + detail::marray_element_type a, + x[i], y[i], a) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(radians, T, degrees, degrees[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step, T edge, T x, edge[i], x[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD( + step, detail::marray_element_type edge, T x, edge, x[i]) +__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(smoothstep, T edge0, T edge1, T x, + edge0[i], edge1[i], x[i]) +__SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD( + smoothstep, detail::marray_element_type edge0, + detail::marray_element_type edge1, T x, edge0, edge1, x[i]) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(sign, T, x, x[i]) + /* --------------- 4.13.4 Integer functions. --------------------------------*/ // ugeninteger abs (geninteger x) template @@ -1724,6 +1908,7 @@ __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(cos) __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(tan) __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(exp) __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(exp2) +__SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(exp10) __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(log) __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(log2) __SYCL_HALF_PRECISION_MATH_FUNCTION_OVERLOAD(log10) From 549d30d51f84d949fa5250bbb5e5a6ea6bb49b1e Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 15 Mar 2023 04:50:17 -0700 Subject: [PATCH 2/8] Fix wrong formatting --- sycl/include/sycl/builtins.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 2b855f316f951..f9021d599cbb9 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -206,7 +206,7 @@ inline __SYCL_ALWAYS_INLINE } __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmax) - __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) +__SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) #undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD From 0a4f0c6cfcd8080580cb6a1d858d82f02c991af0 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 15 Mar 2023 05:06:41 -0700 Subject: [PATCH 3/8] Remove useless undefs --- sycl/include/sycl/builtins.hpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index f9021d599cbb9..c88d39c0b2a4c 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -221,8 +221,6 @@ __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) return res; } -#undef __SYCL_MATH_FUNCTION_2_GENINT_K_OVERLOAD_IMPL - template inline __SYCL_ALWAYS_INLINE std::enable_if_t::value, marray> @@ -234,8 +232,6 @@ inline __SYCL_ALWAYS_INLINE return res; } -#undef __SYCL_MATH_FUNCTION_2_INT_K_OVERLOAD_IMPL - #define __SYCL_MATH_FUNCTION_2_GENINT_Y_OVERLOAD_IMPL(NAME) \ marray res; \ for (size_t i = 0; i < N; i++) { \ From 23cfc1e139caef4be3391dd9eaccf66d4c7d39d4 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 16 Mar 2023 07:42:45 -0700 Subject: [PATCH 4/8] More format --- sycl/include/sycl/builtins.hpp | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index c88d39c0b2a4c..c7616638ee3b2 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -182,8 +182,9 @@ __SYCL_MATH_FUNCTION_2_OVERLOAD(remainder) template inline __SYCL_ALWAYS_INLINE std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray> - powr(marray x, - marray y) __NOEXC{__SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(powr)} + powr(marray x, marray y) __NOEXC { + __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(powr) +} #undef __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL @@ -210,10 +211,10 @@ __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) #undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD - template - inline __SYCL_ALWAYS_INLINE - std::enable_if_t::value, marray> ldexp( - marray x, marray k) __NOEXC { +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + ldexp(marray x, marray k) __NOEXC { marray res; for (size_t i = 0; i < N; i++) { res[i] = __sycl_std::__invoke_ldexp(x[i], k[i]); @@ -272,8 +273,9 @@ inline __SYCL_ALWAYS_INLINE template inline __SYCL_ALWAYS_INLINE std::enable_if_t::value, marray> - rootn(marray x, - int y) __NOEXC{__SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(rootn)} + rootn(marray x, int y) __NOEXC { + __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(rootn) +} #undef __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL From 886dd84d4647e5acdd3fc8d5ec1cb768a835ca84 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 16 Mar 2023 09:08:23 -0700 Subject: [PATCH 5/8] Unify __SYCL_MARRAY_COMMON_FUNCTION* defines --- sycl/include/sycl/builtins.hpp | 48 ++++++++++++++-------------------- 1 file changed, 19 insertions(+), 29 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index c7616638ee3b2..643528f82c2af 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -902,43 +902,33 @@ detail::enable_if_t::value, T> sign(T x) __NOEXC { // TODO: can be optimized in the way math functions are optimized (usage of // vec) -#define __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(NAME, ARGTYPE, ARGVAL, \ - ITEM) \ +#define __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, ...) \ + T res; \ + for (int i = 0; i < T::size(); i++) { \ + res[i] = NAME(__VA_ARGS__); \ + } \ + return res; + +#define __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(NAME, ARG, ...) \ template ::value>> \ - sycl::marray, T::size()> NAME(ARGTYPE ARGVAL) \ - __NOEXC { \ - sycl::marray, T::size()> res; \ - for (int i = 0; i < ARGVAL.size(); i++) { \ - res[i] = NAME(ITEM); \ - } \ - return res; \ + T NAME(ARG) __NOEXC { \ + __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \ } -#define __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(NAME, ARG1, ARG2, ITEM1, \ - ITEM2) \ +#define __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(NAME, ARG1, ARG2, ...) \ template ::value>> \ - sycl::marray, T::size()> NAME(ARG1, ARG2) \ - __NOEXC { \ - sycl::marray, T::size()> res; \ - for (int i = 0; i < x.size(); i++) { \ - res[i] = NAME(ITEM1, ITEM2); \ - } \ - return res; \ + T NAME(ARG1, ARG2) __NOEXC { \ + __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \ } #define __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(NAME, ARG1, ARG2, ARG3, \ - ITEM1, ITEM2, ITEM3) \ + ...) \ template ::value>> \ - sycl::marray, T::size()> NAME(ARG1, ARG2, \ - ARG3) __NOEXC { \ - sycl::marray, T::size()> res; \ - for (int i = 0; i < x.size(); i++) { \ - res[i] = NAME(ITEM1, ITEM2, ITEM3); \ - } \ - return res; \ + T NAME(ARG1, ARG2, ARG3) __NOEXC { \ + __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \ } __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(clamp, T x, T minval, T maxval, @@ -946,7 +936,7 @@ __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(clamp, T x, T minval, T maxval, __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD( clamp, T x, detail::marray_element_type minval, detail::marray_element_type maxval, x[i], minval, maxval) -__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(degrees, T, radians, radians[i]) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(degrees, T radians, radians[i]) __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, T y, x[i], y[i]) __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, detail::marray_element_type y, @@ -960,7 +950,7 @@ __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, T a, x[i], y[i], __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, detail::marray_element_type a, x[i], y[i], a) -__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(radians, T, degrees, degrees[i]) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(radians, T degrees, degrees[i]) __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step, T edge, T x, edge[i], x[i]) __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD( step, detail::marray_element_type edge, T x, edge, x[i]) @@ -969,7 +959,7 @@ __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(smoothstep, T edge0, T edge1, T x, __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD( smoothstep, detail::marray_element_type edge0, detail::marray_element_type edge1, T x, edge0, edge1, x[i]) -__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(sign, T, x, x[i]) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(sign, T x, x[i]) /* --------------- 4.13.4 Integer functions. --------------------------------*/ // ugeninteger abs (geninteger x) From 31fddc1089a3e904ebe01d0986aa5c067186faf4 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Thu, 16 Mar 2023 10:10:49 -0700 Subject: [PATCH 6/8] Add missing undefs --- sycl/include/sycl/builtins.hpp | 37 +++++++++++++++++++++------------- 1 file changed, 23 insertions(+), 14 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 643528f82c2af..24623c4f9ddab 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -916,6 +916,12 @@ detail::enable_if_t::value, T> sign(T x) __NOEXC { __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \ } +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(degrees, T radians, radians[i]) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(radians, T degrees, degrees[i]) +__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(sign, T x, x[i]) + +#undef __SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD + #define __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(NAME, ARG1, ARG2, ...) \ template ::value>> \ @@ -923,6 +929,20 @@ detail::enable_if_t::value, T> sign(T x) __NOEXC { __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL(NAME, __VA_ARGS__) \ } +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(min, T x, T y, x[i], y[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(min, T x, + detail::marray_element_type y, + x[i], y) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, T y, x[i], y[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, + detail::marray_element_type y, + x[i], y) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step, T edge, T x, edge[i], x[i]) +__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD( + step, detail::marray_element_type edge, T x, edge, x[i]) + +#undef __SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD + #define __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(NAME, ARG1, ARG2, ARG3, \ ...) \ template minval, detail::marray_element_type maxval, x[i], minval, maxval) -__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(degrees, T radians, radians[i]) -__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, T y, x[i], y[i]) -__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(max, T x, - detail::marray_element_type y, - x[i], y) -__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(min, T x, T y, x[i], y[i]) -__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(min, T x, - detail::marray_element_type y, - x[i], y) __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, T a, x[i], y[i], a[i]) __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(mix, T x, T y, detail::marray_element_type a, x[i], y[i], a) -__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(radians, T degrees, degrees[i]) -__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD(step, T edge, T x, edge[i], x[i]) -__SYCL_MARRAY_COMMON_FUNCTION_BINOP_OVERLOAD( - step, detail::marray_element_type edge, T x, edge, x[i]) __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD(smoothstep, T edge0, T edge1, T x, edge0[i], edge1[i], x[i]) __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD( smoothstep, detail::marray_element_type edge0, detail::marray_element_type edge1, T x, edge0, edge1, x[i]) -__SYCL_MARRAY_COMMON_FUNCTION_UNOP_OVERLOAD(sign, T x, x[i]) + +#undef __SYCL_MARRAY_COMMON_FUNCTION_TEROP_OVERLOAD +#undef __SYCL_MARRAY_COMMON_FUNCTION_OVERLOAD_IMPL /* --------------- 4.13.4 Integer functions. --------------------------------*/ // ugeninteger abs (geninteger x) From 2a396e01348b743ff75679a81c3d2291ea85c95e Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Fri, 17 Mar 2023 11:16:43 -0700 Subject: [PATCH 7/8] Revert "More format" This reverts commit 23cfc1e139caef4be3391dd9eaccf66d4c7d39d4. --- sycl/include/sycl/builtins.hpp | 18 ++++++++---------- 1 file changed, 8 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 6783a80eb6c48..02a5e6fe00dc7 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -182,9 +182,8 @@ __SYCL_MATH_FUNCTION_2_OVERLOAD(remainder) template inline __SYCL_ALWAYS_INLINE std::enable_if_t<__FAST_MATH_SGENFLOAT(T), marray> - powr(marray x, marray y) __NOEXC { - __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(powr) -} + powr(marray x, + marray y) __NOEXC{__SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL(powr)} #undef __SYCL_MATH_FUNCTION_2_OVERLOAD_IMPL @@ -211,10 +210,10 @@ __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) #undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD -template -inline __SYCL_ALWAYS_INLINE - std::enable_if_t::value, marray> - ldexp(marray x, marray k) __NOEXC { + template + inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> ldexp( + marray x, marray k) __NOEXC { marray res; for (size_t i = 0; i < N; i++) { res[i] = __sycl_std::__invoke_ldexp(x[i], k[i]); @@ -273,9 +272,8 @@ inline __SYCL_ALWAYS_INLINE template inline __SYCL_ALWAYS_INLINE std::enable_if_t::value, marray> - rootn(marray x, int y) __NOEXC { - __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(rootn) -} + rootn(marray x, + int y) __NOEXC{__SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL(rootn)} #undef __SYCL_MATH_FUNCTION_2_INT_Y_OVERLOAD_IMPL From 59d1ce618b7cad1b2001be63b812dc6ce146921f Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Fri, 17 Mar 2023 11:20:53 -0700 Subject: [PATCH 8/8] Turn off format here --- sycl/include/sycl/builtins.hpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 02a5e6fe00dc7..9ff111c14583a 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -206,14 +206,16 @@ inline __SYCL_ALWAYS_INLINE } __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmax) + // clang-format off __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD(fmin) #undef __SYCL_MATH_FUNCTION_2_SGENFLOAT_Y_OVERLOAD - template - inline __SYCL_ALWAYS_INLINE - std::enable_if_t::value, marray> ldexp( - marray x, marray k) __NOEXC { +template +inline __SYCL_ALWAYS_INLINE + std::enable_if_t::value, marray> + ldexp(marray x, marray k) __NOEXC { + // clang-format on marray res; for (size_t i = 0; i < N; i++) { res[i] = __sycl_std::__invoke_ldexp(x[i], k[i]);