From b122ecd190a4f71baaea35971bcd577d400084d5 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:21:35 +0200 Subject: [PATCH 01/13] introduce complex's group algorithms --- include/sycl_ext_complex.hpp | 407 +++++++++++++++++++++++++++++++++++ 1 file changed, 407 insertions(+) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index 661fed2..b8d5ab6 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -391,12 +391,16 @@ struct is_gencomplex std::is_same_v<_Tp, complex> || std::is_same_v<_Tp, complex> || std::is_same_v<_Tp, complex>> {}; +template +inline constexpr bool is_gencomplex_v = is_gencomplex<_Tp>::value; template struct is_genfloat : std::integral_constant || std::is_same_v<_Tp, float> || std::is_same_v<_Tp, sycl::half>> {}; +template +inline constexpr bool is_genfloat_v = is_genfloat<_Tp>::value; template class complex<_Tp, typename std::enable_if::value>::type> { @@ -1257,6 +1261,19 @@ _SYCL_EXT_CPLX_END_NAMESPACE_STD // MARRAY IMPLEMENTATION //////////////////////////////////////////////////////////////////////////////// +_SYCL_EXT_CPLX_BEGIN_NAMESPACE_STD + +template +struct is_mgencomplex : std::false_type {}; + +template +struct is_mgencomplex> : std::integral_constant> {}; + +template +inline constexpr bool is_mgencomplex_v = is_mgencomplex::value; + +_SYCL_EXT_CPLX_END_NAMESPACE_STD + _SYCL_MARRAY_BEGIN_NAMESPACE // marray of complex class specialisation @@ -1681,6 +1698,396 @@ _SYCL_EXT_CPLX_INLINE_VISIBILITY return rtn; } +//////////////////////////////////////////////////////////////////////////////// +// GROUP ALGORITMHS +//////////////////////////////////////////////////////////////////////////////// + +namespace cplex::detail { + +/// Helper traits to check if the type is a sycl::plus +template +struct is_plus : std::integral_constant>> {}; +template +inline constexpr bool is_plus_v = is_plus::value; + +/// Helper traits to check if the type is a sycl:multiplies +template +struct is_multiplies : std::integral_constant>> {}; +template +inline constexpr bool is_multiplies_v = is_multiplies::value; + +/// Wrapper trait to check if the binary operation is supported +template +struct is_binary_op_supported : std::integral_constant::value || detail::is_multiplies::value)> {}; +template +inline constexpr bool is_binary_op_supported_v = is_binary_op_supported::value; + +/// Helper functions to get the init for sycl::plus binary operation when the type is a gencomplex +template +std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v && detail::is_plus_v), T> get_init() { + return T{0, 0}; +} +/// Helper functions to get the init for sycl::multiply binary operation when the type is a gencomplex +template +std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v && detail::is_multiplies::value), T> get_init() { + return T{1, 0}; +} +/// Helper functions to get the init for sycl::plus binary operation when the type is a mgencomplex +template +std::enable_if_t<(is_mgencomplex_v && detail::is_plus::value), T> get_init() { + using Complex = typename T::value_type; + + T result; + std::fill(result.begin(), result.end(), Complex{0, 0}); + return result; +} +/// Helper functions to get the init for sycl::multiply binary operation when the type is a mgencomplex +template +std::enable_if_t<(is_mgencomplex_v && detail::is_multiplies::value), T> get_init() { + using Complex = typename T::value_type; + + T result; + std::fill(result.begin(), result.end(), Complex{1, 0}); + return result; +} + +} + +/* REDUCE_OVER_GROUP'S OVERLOADS */ + +/// Complex specialization +template > && + is_genfloat_v && + is_genfloat_v && + cplex::detail::is_binary_op_supported_v>> +complex reduce_over_group(Group g, complex x, complex init, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + complex result; + + result.real(sycl::reduce_over_group(g, x.real(), init.real(), binary_op)); + result.imag(sycl::reduce_over_group(g, x.imag(), init.imag(), binary_op)); + + return result; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray specialization +template > && + is_gencomplex_v && + is_gencomplex_v && + cplex::detail::is_binary_op_supported_v>> +sycl::marray reduce_over_group(Group g, sycl::marray x, sycl::marray init, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + sycl::marray result; + + sycl::detail::loop([&](size_t s) { + result[s] = reduce_over_group(g, x[s], init[s], binary_op); + }); + + return result; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray and Complex specialization +template > && + (is_gencomplex_v || is_mgencomplex_v) && + cplex::detail::is_binary_op_supported_v>> +T reduce_over_group(Group g, T x, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + auto init = cplex::detail::get_init(); + + return reduce_over_group(g, x, init, binary_op); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/* JOINT_REDUCE'S OVERLOADS */ + +/// Marray and Complex specialization +template > && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v || is_mgencomplex_v) && + cplex::detail::is_binary_op_supported_v>> +T joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + auto partial = cplex::detail::get_init(); + + sycl::detail::for_each(g, first, last, [&](const typename sycl::detail::remove_pointer::type &x) { + partial = binary_op(partial, x); + }); + + return reduce_over_group(g, partial, init, binary_op); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray and Complex specialization +template > && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || is_mgencomplex_v>) && + cplex::detail::is_binary_op_supported_v>> +typename sycl::detail::remove_pointer_t joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + using T = typename sycl::detail::remove_pointer_t; + + auto init = cplex::detail::get_init(); + + return joint_reduce(g, first, last, init, binary_op); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/* INCLUSIVE_SCAN_OVER_GROUP'S OVERLOADS */ + +/// Complex specialization +template > && + is_genfloat_v && + is_genfloat_v && + cplex::detail::is_binary_op_supported_v>> +complex inclusive_scan_over_group(Group g, complex x, BinaryOperation binary_op, complex init) { +#ifdef __SYCL_DEVICE_ONLY__ + complex result; + + result.real(sycl::inclusive_scan_over_group(g, x.real(), binary_op, init.real())); + result.imag(sycl::inclusive_scan_over_group(g, x.imag(), binary_op, init.imag())); + + return result; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray specialization +template > && + is_gencomplex_v && + is_gencomplex_v && + cplex::detail::is_binary_op_supported_v>> +sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, BinaryOperation binary_op, sycl::marray init) { +#ifdef __SYCL_DEVICE_ONLY__ + sycl::marray result; + + sycl::detail::loop([&](size_t s) { + result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); + }); + + return result; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray and Complex specialization +template > && + (is_gencomplex_v || is_mgencomplex_v) && + cplex::detail::is_binary_op_supported_v>> +T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + auto init = cplex::detail::get_init(); + + return inclusive_scan_over_group(g, x, binary_op, init); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/* JOINT_INCLUSIVE_SCAN'S OVERLOADS */ + +/// Complex specialization +template > && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v || is_mgencomplex_v) && + cplex::detail::is_binary_op_supported_v>> +OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) { +#ifdef __SYCL_DEVICE_ONLY__ + std::ptrdiff_t offset = g.get_local_linear_id(); + std::ptrdiff_t stride = g.get_local_linear_range(); + std::ptrdiff_t N = last - first; + + auto roundup = [=](const std::ptrdiff_t &v, const std::ptrdiff_t &divisor) -> std::ptrdiff_t { + return ((v + divisor - 1) / divisor) * divisor; + }; + + typename std::remove_const_t> x; + typename sycl::detail::remove_pointer_t carry = init; + + for (std::ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) { + std::ptrdiff_t i = chunk + offset; + + if (i < N) + x = first[i]; + + typename sycl::detail::remove_pointer_t out = inclusive_scan_over_group(g, x, binary_op, carry); + + if (i < N) + result[i] = out; + + carry = sycl::group_broadcast(g, out, stride - 1); + } + return result + N; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Complex specialization +template > && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v> || is_mgencomplex_v>) && + cplex::detail::is_binary_op_supported_v>> +OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + using T = typename sycl::detail::remove_pointer_t; + + auto init = cplex::detail::get_init(); + + return joint_inclusive_scan(g, first, last, result, binary_op, init); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/* EXCLUSIVE_SCAN_OVER_GROUP'S OVERLOADS */ + +/// Complex specialization +template > && + is_genfloat_v && + is_genfloat_v && + cplex::detail::is_binary_op_supported_v>> +complex exclusive_scan_over_group(Group g, complex x, complex init, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + complex result; + + result.real(sycl::exclusive_scan_over_group(g, x.real(), init.real(), binary_op)); + result.imag(sycl::exclusive_scan_over_group(g, x.imag(), init.imag(), binary_op)); + + return result; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray specialization +template > && + is_gencomplex_v && + is_gencomplex_v && + cplex::detail::is_binary_op_supported_v>> +sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, sycl::marray init, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + sycl::marray result; + + sycl::detail::loop([&](size_t s) { + result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); + }); + + return result; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Marray and Complex specialization +template > && + (is_gencomplex_v || is_mgencomplex_v) && + cplex::detail::is_binary_op_supported_v>> +T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + auto init = cplex::detail::get_init(); + + return exclusive_scan_over_group(g, x, init, binary_op); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/* JOINT_EXCLUSIVE_SCAN'S OVERLOADS */ + +/// Complex specialization +template > && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v || is_mgencomplex_v) && + // + cplex::detail::is_binary_op_supported_v>> +OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + std::ptrdiff_t offset = g.get_local_linear_id(); + std::ptrdiff_t stride = g.get_local_linear_range(); + std::ptrdiff_t N = last - first; + + auto roundup = [=](const std::ptrdiff_t &v, const std::ptrdiff_t &divisor) -> std::ptrdiff_t { + return ((v + divisor - 1) / divisor) * divisor; + }; + + typename std::remove_const_t> x; + typename sycl::detail::remove_pointer_t carry = init; + + for (std::ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) { + std::ptrdiff_t i = chunk + offset; + if (i < N) + x = first[i]; + + typename sycl::detail::remove_pointer_t out = exclusive_scan_over_group(g, x, carry, binary_op); + + if (i < N) + result[i] = out; + + carry = sycl::group_broadcast(g, binary_op(out, x), stride - 1); + } + return result + N; +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + +/// Complex specialization +template > && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || is_mgencomplex_v>) && + (is_gencomplex_v> || is_mgencomplex_v>) && + cplex::detail::is_binary_op_supported_v>> +OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { +#ifdef __SYCL_DEVICE_ONLY__ + using T = typename sycl::detail::remove_pointer_t; + + auto init = cplex::detail::get_init(); + + return joint_exclusive_scan(g, first, last, result, init, binary_op); +#else + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); +#endif +} + _SYCL_EXT_CPLX_END_NAMESPACE_STD #undef _SYCL_MARRAY_BEGIN_NAMESPACE From dd4f27c2c59076c48344593b68ab4d085968c07a Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:23:21 +0200 Subject: [PATCH 02/13] improve test_helper for complex's group algorithms --- tests/test_helper.hpp | 79 +++++++++++++++++++++++++++++++------------ 1 file changed, 58 insertions(+), 21 deletions(-) diff --git a/tests/test_helper.hpp b/tests/test_helper.hpp index 773e980..a5a2ec4 100644 --- a/tests/test_helper.hpp +++ b/tests/test_helper.hpp @@ -1,8 +1,8 @@ #include #include #include -#include -#include +#include + #include #include "sycl_ext_complex.hpp" @@ -67,6 +67,8 @@ template <> struct numeric_limits { } // namespace std #endif +namespace detail { + template inline bool is_nan_or_inf(T x, U y) { return (std::isnan(x.real()) && std::isnan(y.real())) || (std::isnan(x.imag()) && std::isnan(y.imag())) || @@ -74,7 +76,6 @@ template inline bool is_nan_or_inf(T x, U y) { (std::isinf(x.imag()) && std::isinf(y.imag())); } -namespace detail { template bool almost_equal(T x, T y, int ulp) { if (std::isnan(x) && std::isnan(y)) return true; @@ -87,11 +88,12 @@ template bool almost_equal(T x, T y, int ulp) { } template -bool almost_equal(sycl::ext::cplx::complex x, std::complex y, int ulp) { - auto diff = std::abs((std::complex)x - y); +bool almost_equal(std::complex output, std::complex reference, int ulp) { + auto diff = std::abs(output - reference); return diff <= std::numeric_limits::epsilon() * - std::abs((std::complex)x + y) * ulp || - diff < std::numeric_limits::min() || is_nan_or_inf(x, y); + std::abs(output + reference) * ulp || + diff < std::numeric_limits::min() || + is_nan_or_inf(output, reference); } // Helpers for testing half @@ -106,6 +108,7 @@ inline std::complex trunc_float(std::complex c) { auto c_sycl_half = static_cast>(c); return sycl_half_to_float(c_sycl_half); } + } // namespace detail // Helper for initializing std::complex values for tests only needed because @@ -164,34 +167,68 @@ auto constexpr convert_marray(sycl::marray, NumElements> c) { // Helpers for comparing SyclCPLX and standard c++ results +namespace detail { + +template +struct is_a_complex + : std::integral_constant || + std::is_same_v, T> || + std::is_same_v, T> || + std::is_same_v, T>> {}; +template +inline constexpr bool is_a_complex_v = is_a_complex::value; + +} // namespace detail + +/// Specialization for double, float, sycl::half template -void check_results(sycl::ext::cplx::complex output, - std::complex reference, int tol_multiplier = 1) { +typename std::enable_if_t, void> +check_results(T output, T reference, int tol_multiplier = 1) { CHECK(detail::almost_equal(output, reference, tol_multiplier * SYCL_CPLX_TOL_ULP)); } -template -void check_results(T output, T reference, int tol_multiplier = 1) { - CHECK(detail::almost_equal(output, reference, +/// Specialization for sycl::complex and std::complex +template +typename std::enable_if_t< + detail::is_a_complex_v && detail::is_a_complex_v, void> +check_results(LHS output, RHS reference, int tol_multiplier = 1) { + using T1 = typename LHS::value_type; + using T2 = typename RHS::value_type; + + if (!std::is_same_v) { + FAIL("check_results can be called with sycl::complex and/or std::complex " + "but with the same value_type"); + } + + CHECK(detail::almost_equal(static_cast>(output), + static_cast>(reference), tol_multiplier * SYCL_CPLX_TOL_ULP)); } -template -void check_results( - sycl::marray, NumElements> output, - sycl::marray, NumElements> reference, - int tol_multiplier = 1) { - for (std::size_t i = 0; i < NumElements; ++i) { +/// Specialization for sycl::marray +template +void check_results(sycl::marray output, + sycl::marray reference, + int tol_multiplier = 1) { + if (NumElementsLHS != NumElementsRHS) { + FAIL("check_results can be called with sycl::marray but with the same " + "NumElement"); + } + + for (std::size_t i = 0; i < NumElementsLHS; ++i) { check_results(output[i], reference[i], tol_multiplier); } } +/// Specialization for std::array template -void check_results(sycl::marray output, - sycl::marray reference, +void check_results(std::array output, + std::array reference, int tol_multiplier = 1) { for (std::size_t i = 0; i < NumElements; ++i) { - check_results(output[i], reference[i], tol_multiplier); + check_results(output[i], reference[i], tol_multiplier); } } From 64c05a9388c29c60fda5af62047ccff4397d1a6e Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:24:27 +0200 Subject: [PATCH 03/13] add and group all traits testing together --- tests/test_gencomplex.cpp | 17 ----- tests/test_traits.cpp | 130 ++++++++++++++++++++++++++++++++++++++ 2 files changed, 130 insertions(+), 17 deletions(-) delete mode 100644 tests/test_gencomplex.cpp create mode 100644 tests/test_traits.cpp diff --git a/tests/test_gencomplex.cpp b/tests/test_gencomplex.cpp deleted file mode 100644 index 64b0bd6..0000000 --- a/tests/test_gencomplex.cpp +++ /dev/null @@ -1,17 +0,0 @@ -#include "test_helper.hpp" - -using namespace sycl::ext::cplx; - -// Check is_gencomplex -TEST_CASE("Test is_gencomplex", "[gencomplex]") { - static_assert(is_gencomplex>::value == true); - static_assert(is_gencomplex>::value == true); - static_assert(is_gencomplex>::value == true); - - static_assert(is_gencomplex>::value == false); - static_assert(is_gencomplex>::value == false); - static_assert(is_gencomplex>::value == false); - static_assert(is_gencomplex>::value == false); - static_assert(is_gencomplex>::value == false); - static_assert(is_gencomplex>::value == false); -} diff --git a/tests/test_traits.cpp b/tests/test_traits.cpp new file mode 100644 index 0000000..7473862 --- /dev/null +++ b/tests/test_traits.cpp @@ -0,0 +1,130 @@ +#include "test_helper.hpp" + +using namespace sycl::ext::cplx; + +// Check is_gencomplex +TEST_CASE("Test is_gencomplex", "[traits]") { + static_assert(is_gencomplex>::value == true); + static_assert(is_gencomplex>::value == true); + static_assert(is_gencomplex>::value == true); + + static_assert(is_gencomplex_v> == true); + static_assert(is_gencomplex_v> == true); + static_assert(is_gencomplex_v> == true); + + static_assert(is_gencomplex>::value == false); + static_assert(is_gencomplex>::value == false); + static_assert(is_gencomplex>::value == false); + static_assert(is_gencomplex>::value == false); + static_assert(is_gencomplex>::value == false); + static_assert(is_gencomplex>::value == false); + + static_assert(is_gencomplex_v> == false); + static_assert(is_gencomplex_v> == false); + static_assert(is_gencomplex_v> == false); + static_assert(is_gencomplex_v> == false); + static_assert(is_gencomplex_v> == false); + static_assert(is_gencomplex_v> == false); +} + +// Check is_genfloat +TEST_CASE("Test is_genfloat", "[traits]") { + static_assert(is_genfloat::value == true); + static_assert(is_genfloat::value == true); + static_assert(is_genfloat::value == true); + + static_assert(is_genfloat_v == true); + static_assert(is_genfloat_v == true); + static_assert(is_genfloat_v == true); + + static_assert(is_genfloat::value == false); + static_assert(is_genfloat::value == false); + static_assert(is_genfloat::value == false); + static_assert(is_genfloat::value == false); + static_assert(is_genfloat::value == false); + static_assert(is_genfloat::value == false); + + static_assert(is_genfloat_v == false); + static_assert(is_genfloat_v == false); + static_assert(is_genfloat_v == false); + static_assert(is_genfloat_v == false); + static_assert(is_genfloat_v == false); + static_assert(is_genfloat_v == false); +} + +// Check is_mgencomplex +TEST_CASE("Test is_mgencomplex", "[traits]") { + static_assert(is_mgencomplex, 42>>::value == + true); + static_assert(is_mgencomplex, 42>>::value == + true); + static_assert(is_mgencomplex, 42>>::value == + true); + + static_assert(is_mgencomplex_v, 42>> == true); + static_assert(is_mgencomplex_v, 42>> == true); + static_assert(is_mgencomplex_v, 42>> == + true); + + static_assert(is_mgencomplex, 42>>::value == + false); + static_assert(is_mgencomplex, 42>>::value == + false); + static_assert(is_mgencomplex, 42>>::value == false); + static_assert( + is_mgencomplex, 42>>::value == + false); + static_assert( + is_mgencomplex, 42>>::value == false); + static_assert( + is_mgencomplex, 42>>::value == false); + + static_assert(is_mgencomplex_v, 42>> == + false); + static_assert(is_mgencomplex_v, 42>> == false); + static_assert(is_mgencomplex_v, 42>> == false); + static_assert( + is_mgencomplex_v, 42>> == false); + static_assert(is_mgencomplex_v, 42>> == + false); + static_assert(is_mgencomplex_v, 42>> == + false); +} + +// Check is_plus +TEST_CASE("Test is_plus", "[traits]") { + static_assert(cplex::detail::is_plus_v> == true); + + static_assert(cplex::detail::is_plus_v> == false); + static_assert(cplex::detail::is_plus_v> == false); + static_assert(cplex::detail::is_plus_v> == false); + static_assert(cplex::detail::is_plus_v> == false); + static_assert(cplex::detail::is_plus_v> == false); +} + +// Check is_multiplies +TEST_CASE("Test is_multiplies", "[traits]") { + static_assert(cplex::detail::is_multiplies_v> == true); + + static_assert(cplex::detail::is_multiplies_v> == false); + static_assert(cplex::detail::is_multiplies_v> == false); + static_assert(cplex::detail::is_multiplies_v> == false); + static_assert(cplex::detail::is_multiplies_v> == false); + static_assert(cplex::detail::is_multiplies_v> == false); +} + +// Check is_binary_op_supported +TEST_CASE("Test is_binary_op_supported", "[traits]") { + static_assert(cplex::detail::is_binary_op_supported_v> == true); + static_assert(cplex::detail::is_binary_op_supported_v> == + true); + + static_assert(cplex::detail::is_binary_op_supported_v> == + false); + static_assert(cplex::detail::is_binary_op_supported_v> == + false); + static_assert(cplex::detail::is_binary_op_supported_v> == + false); + static_assert(cplex::detail::is_binary_op_supported_v> == + false); +} From 3795a0ab5edd8bdc555426d6f8e351841b957902 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:25:12 +0200 Subject: [PATCH 04/13] add conversion test to support modification made to the test helper --- tests/test_conversion.cpp | 70 +++++++++++++++++++++++++++++++++++++++ 1 file changed, 70 insertions(+) create mode 100644 tests/test_conversion.cpp diff --git a/tests/test_conversion.cpp b/tests/test_conversion.cpp new file mode 100644 index 0000000..2e5f19e --- /dev/null +++ b/tests/test_conversion.cpp @@ -0,0 +1,70 @@ +#include "test_helper.hpp" + +#include +#include + +// Check conversion sycl:complex to std::complex +TEMPLATE_TEST_CASE("Test sycl::complex to std::complex conversion", + "[conversion]", double, float, sycl::half) { + using T = TestType; + + auto arr = std::array, 8>{ + sycl::ext::cplx::complex{0, 0}, sycl::ext::cplx::complex{-0, 0}, + sycl::ext::cplx::complex{0, -0}, sycl::ext::cplx::complex{-0, -0}, + + sycl::ext::cplx::complex{1, 1}, sycl::ext::cplx::complex{-1, 1}, + sycl::ext::cplx::complex{1, -1}, sycl::ext::cplx::complex{-1, -1}, + }; + + for (const auto &lhs : arr) { + const auto rhs = static_cast>(lhs); + + assert(lhs.real() == rhs.real() && + "sycl::complex differs from std::complex after conversion"); + assert(lhs.imag() == rhs.imag() && + "sycl::complex differs from std::complex after conversion"); + } +} + +// Check edge-cases conversion sycl:complex to std::complex +TEMPLATE_TEST_CASE("Test edge-cases sycl::complex to std::complex conversion", + "[conversion]", double, float, sycl::half) { + using T = TestType; + + { + auto lhs = sycl::ext::cplx::complex{inf_val, inf_val}; + auto rhs = static_cast>(lhs); + + assert(std::isinf(lhs.real()) && std::isinf(rhs.real()) && + "sycl::complex differs from std::complex after conversion"); + assert(std::isinf(lhs.imag()) && std::isinf(rhs.imag()) && + "sycl::complex differs from std::complex after conversion"); + } + { + auto lhs = sycl::ext::cplx::complex{inf_val, nan_val}; + auto rhs = static_cast>(lhs); + + assert(std::isinf(lhs.real()) && std::isinf(rhs.real()) && + "sycl::complex differs from std::complex after conversion"); + assert(std::isnan(lhs.imag()) && std::isnan(rhs.imag()) && + "sycl::complex differs from std::complex after conversion"); + } + { + auto lhs = sycl::ext::cplx::complex{nan_val, inf_val}; + auto rhs = static_cast>(lhs); + + assert(std::isnan(lhs.real()) && std::isnan(rhs.real()) && + "sycl::complex differs from std::complex after conversion"); + assert(std::isinf(lhs.imag()) && std::isinf(rhs.imag()) && + "sycl::complex differs from std::complex after conversion"); + } + { + auto lhs = sycl::ext::cplx::complex{nan_val, nan_val}; + auto rhs = static_cast>(lhs); + + assert(std::isnan(lhs.real()) && std::isnan(rhs.real()) && + "sycl::complex differs from std::complex after conversion"); + assert(std::isnan(lhs.imag()) && std::isnan(rhs.imag()) && + "sycl::complex differs from std::complex after conversion"); + } +} From 690e0a56b6d9ffd7e64e9136163c8b5d7b83a7fe Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:26:05 +0200 Subject: [PATCH 05/13] add tests for complex's group algorithms --- tests/exclusive_scan_over_group.cpp | 191 ++++++++++++++++++++++++++++ tests/inclusive_scan_over_group.cpp | 190 +++++++++++++++++++++++++++ tests/joint_exclusive_scan.cpp | 190 +++++++++++++++++++++++++++ tests/joint_inclusive_scan.cpp | 188 +++++++++++++++++++++++++++ tests/joint_reduce.cpp | 178 ++++++++++++++++++++++++++ tests/reduce_over_group.cpp | 179 ++++++++++++++++++++++++++ 6 files changed, 1116 insertions(+) create mode 100644 tests/exclusive_scan_over_group.cpp create mode 100644 tests/inclusive_scan_over_group.cpp create mode 100644 tests/joint_exclusive_scan.cpp create mode 100644 tests/joint_inclusive_scan.cpp create mode 100644 tests/joint_reduce.cpp create mode 100644 tests/reduce_over_group.cpp diff --git a/tests/exclusive_scan_over_group.cpp b/tests/exclusive_scan_over_group.cpp new file mode 100644 index 0000000..469c1ef --- /dev/null +++ b/tests/exclusive_scan_over_group.cpp @@ -0,0 +1,191 @@ +#include +#include + +#include "test_helper.hpp" + +//////////////////////////////////////////////////////////////////////////////// +// UTILITY FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// + +template +void test_exclusive_scan_over_group(sycl::queue q, T input, + BinaryOperation binary_op) { + using V = typename T::value_type; + + constexpr size_t N = input.size(); + + auto init = sycl::ext::cplx::cplex::detail::get_init(); + + auto *in = sycl::malloc_shared(N, q); + auto *output_with_init = sycl::malloc_shared(N, q); + auto *output_without_init = sycl::malloc_shared(N, q); + + for (std::size_t i = 0; i < N; i++) { + in[i] = input[i]; + } + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) { + auto gid = it.get_global_id(0); + auto lid = it.get_local_id(0); + auto g = it.get_group(); + + output_with_init[lid] = sycl::ext::cplx::exclusive_scan_over_group( + g, in[gid], init, binary_op); + output_without_init[lid] = + sycl::ext::cplx::exclusive_scan_over_group(g, in[gid], binary_op); + }); + }); + + q.wait(); + + std::array expected; + std::exclusive_scan(input.begin(), input.end(), expected.begin(), init, + binary_op); + + std::array result; + for (std::size_t i = 0; i < N; i++) { + result[i] = output_with_init[i]; + } + + check_results(result, expected); + + for (std::size_t i = 0; i < N; i++) { + result[i] = output_without_init[i]; + } + + check_results(result, expected); + + sycl::free(in, q); + sycl::free(output_with_init, q); + sycl::free(output_without_init, q); +} + +//////////////////////////////////////////////////////////////////////////////// +// COMPLEX TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test complex exclusive_scan_over_group", "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}}, + // Random value test + Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + // Repeated value test + Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + // Negative value test + Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + // Large value test + Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + // Small value test + Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + // Edge case value test + Array{Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_exclusive_scan_over_group(q, test_cases, binary_op); + } +} + +//////////////////////////////////////////////////////////////////////////////// +// MARRAY TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test marray exclusive_scan_over_group", + "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Marray = typename sycl::marray; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}}, + Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}}, + Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}}, + Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}}, + // Random value test + Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}}, + // Repeated value test + Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}}, + // Negative value test + Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}}, + // Large value test + Array{ + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}}, + // Small value test + Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}}, + // Edge case value test + Array{ + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{Complex{nan_val, nan_val}, + Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, + Complex{inf_val, nan_val}}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_exclusive_scan_over_group(q, test_cases, binary_op); + } +} diff --git a/tests/inclusive_scan_over_group.cpp b/tests/inclusive_scan_over_group.cpp new file mode 100644 index 0000000..e84fe4a --- /dev/null +++ b/tests/inclusive_scan_over_group.cpp @@ -0,0 +1,190 @@ +#include +#include + +#include "test_helper.hpp" + +//////////////////////////////////////////////////////////////////////////////// +// UTILITY FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// + +template +void test_inclusive_scan_over_group(sycl::queue q, T input, + BinaryOperation binary_op) { + using V = typename T::value_type; + + constexpr size_t N = input.size(); + + auto init = sycl::ext::cplx::cplex::detail::get_init(); + + auto *in = sycl::malloc_shared(N, q); + auto *output_with_init = sycl::malloc_shared(N, q); + auto *output_without_init = sycl::malloc_shared(N, q); + + for (std::size_t i = 0; i < N; i++) { + in[i] = input[i]; + } + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) { + auto gid = it.get_global_id(0); + auto lid = it.get_local_id(0); + auto g = it.get_group(); + + output_with_init[lid] = sycl::ext::cplx::inclusive_scan_over_group( + g, in[gid], binary_op, init); + output_without_init[lid] = + sycl::ext::cplx::inclusive_scan_over_group(g, in[gid], binary_op); + }); + }); + + q.wait(); + + std::array expected; + std::inclusive_scan(input.begin(), input.end(), expected.begin(), binary_op); + + std::array result; + for (std::size_t i = 0; i < N; i++) { + result[i] = output_with_init[i]; + } + + check_results(result, expected); + + for (std::size_t i = 0; i < N; i++) { + result[i] = output_without_init[i]; + } + + check_results(result, expected); + + sycl::free(in, q); + sycl::free(output_with_init, q); + sycl::free(output_without_init, q); +} + +//////////////////////////////////////////////////////////////////////////////// +// COMPLEX TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test complex inclusive_scan_over_group", "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}}, + // Random value test + Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + // Repeated value test + Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + // Negative value test + Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + // Large value test + Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + // Small value test + Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + // Edge case value test + Array{Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_inclusive_scan_over_group(q, test_cases, binary_op); + } +} + +//////////////////////////////////////////////////////////////////////////////// +// MARRAY TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test marray inclusive_scan_over_group", + "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Marray = typename sycl::marray; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}}, + Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}}, + Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}}, + Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}}, + // Random value test + Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}}, + // Repeated value test + Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}}, + // Negative value test + Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}}, + // Large value test + Array{ + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}}, + // Small value test + Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}}, + // Edge case value test + Array{ + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{Complex{nan_val, nan_val}, + Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, + Complex{inf_val, nan_val}}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_inclusive_scan_over_group(q, test_cases, binary_op); + } +} diff --git a/tests/joint_exclusive_scan.cpp b/tests/joint_exclusive_scan.cpp new file mode 100644 index 0000000..dd3584d --- /dev/null +++ b/tests/joint_exclusive_scan.cpp @@ -0,0 +1,190 @@ +#include +#include + +#include "test_helper.hpp" + +//////////////////////////////////////////////////////////////////////////////// +// UTILITY FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// + +template +void test_joint_exclusive_scan(sycl::queue q, T input, + BinaryOperation binary_op) { + using V = typename T::value_type; + + constexpr size_t N = input.size(); + + auto init = sycl::ext::cplx::cplex::detail::get_init(); + + auto *in = sycl::malloc_shared(N, q); + auto *output_with_init = sycl::malloc_shared(N, q); + auto *output_without_init = sycl::malloc_shared(N, q); + + for (std::size_t i = 0; i < N; i++) { + in[i] = input[i]; + } + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) { + auto gid = it.get_global_id(0); + auto lid = it.get_local_id(0); + auto g = it.get_group(); + + sycl::ext::cplx::joint_exclusive_scan(g, in, in + N, output_with_init, + init, binary_op); + sycl::ext::cplx::joint_exclusive_scan(g, in, in + N, output_without_init, + binary_op); + }); + }); + + q.wait(); + + std::array expected; + std::exclusive_scan(input.begin(), input.end(), expected.begin(), init, + binary_op); + + std::array result; + for (std::size_t i = 0; i < N; i++) { + result[i] = output_with_init[i]; + } + + check_results(result, expected); + + for (std::size_t i = 0; i < N; i++) { + result[i] = output_without_init[i]; + } + + check_results(result, expected); + + sycl::free(in, q); + sycl::free(output_with_init, q); + sycl::free(output_without_init, q); +} + +//////////////////////////////////////////////////////////////////////////////// +// COMPLEX TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test complex joint_exclusive_scan", "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}}, + // Random value test + Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + // Repeated value test + Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + // Negative value test + Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + // Large value test + Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + // Small value test + Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + // Edge case value test + Array{Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_joint_exclusive_scan(q, test_cases, binary_op); + } +} + +//////////////////////////////////////////////////////////////////////////////// +// MARRAY TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test marray joint_exclusive_scan", "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Marray = typename sycl::marray; + using Array = typename std::array; + + sycl::queue q; + + const auto binary_op = BinaryOperation{}; + + const auto test_cases = GENERATE( + // Basic value test + Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}}, + Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}}, + Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}}, + Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}}, + // Random value test + Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}}, + // Repeated value test + Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}}, + // Negative value test + Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}}, + // Large value test + Array{ + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}}, + // Small value test + Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}}, + // Edge case value test + Array{ + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{Complex{nan_val, nan_val}, + Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, + Complex{inf_val, nan_val}}}); + if (is_type_supported(q)) { + test_joint_exclusive_scan(q, test_cases, binary_op); + } +} diff --git a/tests/joint_inclusive_scan.cpp b/tests/joint_inclusive_scan.cpp new file mode 100644 index 0000000..1577fc4 --- /dev/null +++ b/tests/joint_inclusive_scan.cpp @@ -0,0 +1,188 @@ +#include +#include + +#include "test_helper.hpp" + +//////////////////////////////////////////////////////////////////////////////// +// UTILITY FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// + +template +void test_joint_inclusive_scan(sycl::queue q, T input, + BinaryOperation binary_op) { + using V = typename T::value_type; + + constexpr size_t N = input.size(); + + auto init = sycl::ext::cplx::cplex::detail::get_init(); + + auto *in = sycl::malloc_shared(N, q); + auto *output_with_init = sycl::malloc_shared(N, q); + auto *output_without_init = sycl::malloc_shared(N, q); + + for (std::size_t i = 0; i < N; i++) { + in[i] = input[i]; + } + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) { + auto gid = it.get_global_id(0); + auto lid = it.get_local_id(0); + auto g = it.get_group(); + + sycl::ext::cplx::joint_inclusive_scan(g, in, in + N, output_with_init, + binary_op, init); + sycl::ext::cplx::joint_inclusive_scan(g, in, in + N, output_without_init, + binary_op); + }); + }); + + q.wait(); + + std::array expected; + std::inclusive_scan(input.begin(), input.end(), expected.begin(), binary_op); + + std::array result; + for (std::size_t i = 0; i < N; i++) { + result[i] = output_with_init[i]; + } + + check_results(result, expected); + + for (std::size_t i = 0; i < N; i++) { + result[i] = output_without_init[i]; + } + + check_results(result, expected); + + sycl::free(in, q); + sycl::free(output_with_init, q); + sycl::free(output_without_init, q); +} + +//////////////////////////////////////////////////////////////////////////////// +// COMPLEX TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test complex joint_inclusive_scan", "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}}, + // Random value test + Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + // Repeated value test + Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + // Negative value test + Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + // Large value test + Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + // Small value test + Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + // Edge case value test + Array{Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_joint_inclusive_scan(q, test_cases, binary_op); + } +} + +//////////////////////////////////////////////////////////////////////////////// +// MARRAY TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test marray joint_inclusive_scan", "[scan]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Marray = typename sycl::marray; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}}, + Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}}, + Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}}, + Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}}, + // Random value test + Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}}, + // Repeated value test + Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}}, + // Negative value test + Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}}, + // Large value test + Array{ + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}}, + // Small value test + Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}}, + // Edge case value test + Array{ + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{Complex{nan_val, nan_val}, + Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, + Complex{inf_val, nan_val}}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_joint_inclusive_scan(q, test_cases, binary_op); + } +} diff --git a/tests/joint_reduce.cpp b/tests/joint_reduce.cpp new file mode 100644 index 0000000..9e22158 --- /dev/null +++ b/tests/joint_reduce.cpp @@ -0,0 +1,178 @@ +#include +#include + +#include "test_helper.hpp" + +//////////////////////////////////////////////////////////////////////////////// +// UTILITY FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// + +template +void test_joint_reduce(sycl::queue q, T input, BinaryOperation binary_op) { + using V = typename T::value_type; + + constexpr size_t N = input.size(); + + auto init = sycl::ext::cplx::cplex::detail::get_init(); + + auto *in = sycl::malloc_shared(N, q); + auto *output_with_init = sycl::malloc_shared(1, q); + auto *output_without_init = sycl::malloc_shared(1, q); + + for (std::size_t i = 0; i < N; i++) { + in[i] = input[i]; + } + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) { + auto gid = it.get_global_id(0); + auto lid = it.get_local_id(0); + auto g = it.get_group(); + + *output_with_init = + sycl::ext::cplx::joint_reduce(g, in, in + N, init, binary_op); + *output_without_init = + sycl::ext::cplx::joint_reduce(g, in, in + N, binary_op); + }); + }); + + q.wait(); + + const auto expected = + std::reduce(input.begin(), input.end(), init, binary_op); + + check_results(*output_with_init, expected); + check_results(*output_without_init, expected); + + sycl::free(in, q); + sycl::free(output_with_init, q); + sycl::free(output_without_init, q); +} + +//////////////////////////////////////////////////////////////////////////////// +// COMPLEX TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test complex joint_reduce", "[reduction]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}}, + // Random value test + Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + // Repeated value test + Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + // Negative value test + Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + // Large value test + Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + // Small value test + Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + // Edge case value test + Array{Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_joint_reduce(q, test_cases, binary_op); + } +} + +//////////////////////////////////////////////////////////////////////////////// +// MARRAY TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test marray joint_reduce", "[reduction]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Marray = typename sycl::marray; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}}, + Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}}, + Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}}, + Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}}, + // Random value test + Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}}, + // Repeated value test + Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}}, + // Negative value test + Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}}, + // Large value test + Array{ + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}}, + // Small value test + Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}}, + // Edge case value test + Array{ + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{Complex{nan_val, nan_val}, + Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, + Complex{inf_val, nan_val}}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_joint_reduce(q, test_cases, binary_op); + } +} diff --git a/tests/reduce_over_group.cpp b/tests/reduce_over_group.cpp new file mode 100644 index 0000000..2728c1c --- /dev/null +++ b/tests/reduce_over_group.cpp @@ -0,0 +1,179 @@ +#include +#include + +#include "test_helper.hpp" + +//////////////////////////////////////////////////////////////////////////////// +// UTILITY FUNCTIONS +//////////////////////////////////////////////////////////////////////////////// + +template +void test_reduce_over_group(sycl::queue q, T input, BinaryOperation binary_op) { + using V = typename T::value_type; + + constexpr size_t N = input.size(); + + const auto init = + sycl::ext::cplx::cplex::detail::get_init(); + + auto *in = sycl::malloc_shared(N, q); + auto *output_with_init = sycl::malloc_shared(1, q); + auto *output_without_init = sycl::malloc_shared(1, q); + + for (std::size_t i = 0; i < N; i++) { + in[i] = input[i]; + } + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) { + auto gid = it.get_global_id(0); + auto lid = it.get_local_id(0); + auto g = it.get_group(); + + *output_with_init = + sycl::ext::cplx::reduce_over_group(g, in[gid], init, binary_op); + *output_without_init = + sycl::ext::cplx::reduce_over_group(g, in[gid], binary_op); + }); + }); + + q.wait(); + + const auto expected = + std::reduce(input.begin(), input.end(), init, binary_op); + + check_results(*output_with_init, expected); + check_results(*output_without_init, expected); + + sycl::free(in, q); + sycl::free(output_with_init, q); + sycl::free(output_without_init, q); +} + +//////////////////////////////////////////////////////////////////////////////// +// COMPLEX TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test complex reduce_over_group", "[reduction]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}}, + // Random value test + Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + // Repeated value test + Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + // Negative value test + Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + // Large value test + Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + // Small value test + Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + // Edge case value test + Array{Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_reduce_over_group(q, test_cases, binary_op); + } +} + +//////////////////////////////////////////////////////////////////////////////// +// MARRAY TESTS +//////////////////////////////////////////////////////////////////////////////// + +TEMPLATE_TEST_CASE_SIG("Test marray reduce_over_group", "[reduction]", + ((typename T, std::size_t N, typename BinaryOperation), + T, N, BinaryOperation), + (double, 4, sycl::plus<>), (float, 4, sycl::plus<>), + (sycl::half, 4, sycl::plus<>)) { + + using Complex = typename sycl::ext::cplx::complex; + using Marray = typename sycl::marray; + using Array = typename std::array; + + sycl::queue q; + + const auto test_cases = GENERATE( + // Basic value test + Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}}, + Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}}, + Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}}, + Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}}, + // Random value test + Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}, + Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8}, + Complex{3.7, 3.7}}}, + // Repeated value test + Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}, + Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}}, + // Negative value test + Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}, + Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2}, + Complex{0, 0}}}, + // Large value test + Array{ + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}, + Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0}, + Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}}, + // Small value test + Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}, + Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002}, + Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}}, + // Edge case value test + Array{ + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{ + Complex{nan_val, nan_val}, Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, Complex{inf_val, nan_val}}, + Marray{Complex{nan_val, nan_val}, + Complex{inf_val, inf_val}, + Complex{nan_val, inf_val}, + Complex{inf_val, nan_val}}}); + const auto binary_op = BinaryOperation{}; + + if (is_type_supported(q)) { + test_reduce_over_group(q, test_cases, binary_op); + } +} From 1655ac9b176a55bfaac26da58ccf308d283c44b1 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:26:38 +0200 Subject: [PATCH 06/13] formatting --- tests/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 86f9467..bbd0f54 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -3,7 +3,7 @@ file(GLOB test_cases CONFIGURE_DEPENDS "*.cpp") foreach(test_file IN LISTS test_cases) if(EXISTS "${test_file}") get_filename_component(exe_name "${test_file}" NAME_WE) - + add_executable(${exe_name} ${test_file}) target_include_directories(${exe_name} PUBLIC ../include/) target_link_libraries(${exe_name} PRIVATE From a905ea32c899ea462af03ea40a9eab110ba15314 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:27:24 +0200 Subject: [PATCH 07/13] increase tol for device's asin --- tests/asin_complex.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/asin_complex.cpp b/tests/asin_complex.cpp index 6c161b0..97902b9 100644 --- a/tests/asin_complex.cpp +++ b/tests/asin_complex.cpp @@ -47,7 +47,7 @@ TEMPLATE_TEST_CASE("Test complex asin", "[asin]", double, float, sycl::half) { } Q.wait(); - check_results(cplx_out[0], std_out, /*tol_multiplier*/ 5); + check_results(cplx_out[0], std_out, /*tol_multiplier*/ 6); } // Check cplx::complex output from host From 77ab067a7ea121dc063c7c7f1ba91ff5f6849aba Mon Sep 17 00:00:00 2001 From: jle-quel Date: Thu, 1 Jun 2023 09:42:01 +0200 Subject: [PATCH 08/13] apply same version of clang-format than the CI --- include/sycl_ext_complex.hpp | 391 ++++++++++++++++++++++------------- 1 file changed, 246 insertions(+), 145 deletions(-) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index b8d5ab6..22fc617 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -310,8 +310,8 @@ template <> struct __numeric_type { }; template ::value &&__numeric_type<_A2>::value - &&__numeric_type<_A3>::value> + bool = __numeric_type<_A1>::value && __numeric_type<_A2>::value && + __numeric_type<_A3>::value> class __promote_imp { public: static const bool value = false; @@ -1263,11 +1263,11 @@ _SYCL_EXT_CPLX_END_NAMESPACE_STD _SYCL_EXT_CPLX_BEGIN_NAMESPACE_STD -template -struct is_mgencomplex : std::false_type {}; +template struct is_mgencomplex : std::false_type {}; template -struct is_mgencomplex> : std::integral_constant> {}; +struct is_mgencomplex> + : std::integral_constant> {}; template inline constexpr bool is_mgencomplex_v = is_mgencomplex::value; @@ -1706,44 +1706,68 @@ namespace cplex::detail { /// Helper traits to check if the type is a sycl::plus template -struct is_plus : std::integral_constant>> {}; +struct is_plus + : std::integral_constant>> { +}; template inline constexpr bool is_plus_v = is_plus::value; /// Helper traits to check if the type is a sycl:multiplies template -struct is_multiplies : std::integral_constant>> {}; +struct is_multiplies + : std::integral_constant< + bool, std::is_same_v>> {}; template inline constexpr bool is_multiplies_v = is_multiplies::value; /// Wrapper trait to check if the binary operation is supported template -struct is_binary_op_supported : std::integral_constant::value || detail::is_multiplies::value)> {}; +struct is_binary_op_supported + : std::integral_constant::value || + detail::is_multiplies::value)> { +}; template -inline constexpr bool is_binary_op_supported_v = is_binary_op_supported::value; +inline constexpr bool is_binary_op_supported_v = + is_binary_op_supported::value; -/// Helper functions to get the init for sycl::plus binary operation when the type is a gencomplex +/// Helper functions to get the init for sycl::plus binary operation when the +/// type is a gencomplex template -std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v && detail::is_plus_v), T> get_init() { +std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v && + detail::is_plus_v), + T> +get_init() { return T{0, 0}; } -/// Helper functions to get the init for sycl::multiply binary operation when the type is a gencomplex +/// Helper functions to get the init for sycl::multiply binary operation when +/// the type is a gencomplex template -std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v && detail::is_multiplies::value), T> get_init() { +std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v && + detail::is_multiplies::value), + T> +get_init() { return T{1, 0}; } -/// Helper functions to get the init for sycl::plus binary operation when the type is a mgencomplex +/// Helper functions to get the init for sycl::plus binary operation when the +/// type is a mgencomplex template -std::enable_if_t<(is_mgencomplex_v && detail::is_plus::value), T> get_init() { +std::enable_if_t< + (is_mgencomplex_v && detail::is_plus::value), T> +get_init() { using Complex = typename T::value_type; T result; std::fill(result.begin(), result.end(), Complex{0, 0}); return result; } -/// Helper functions to get the init for sycl::multiply binary operation when the type is a mgencomplex +/// Helper functions to get the init for sycl::multiply binary operation when +/// the type is a mgencomplex template -std::enable_if_t<(is_mgencomplex_v && detail::is_multiplies::value), T> get_init() { +std::enable_if_t< + (is_mgencomplex_v && detail::is_multiplies::value), T> +get_init() { using Complex = typename T::value_type; T result; @@ -1751,17 +1775,18 @@ std::enable_if_t<(is_mgencomplex_v && detail::is_multiplies: return result; } -} +} // namespace cplex::detail /* REDUCE_OVER_GROUP'S OVERLOADS */ /// Complex specialization -template > && - is_genfloat_v && - is_genfloat_v && - cplex::detail::is_binary_op_supported_v>> -complex reduce_over_group(Group g, complex x, complex init, BinaryOperation binary_op) { +template > && is_genfloat_v && + is_genfloat_v && + cplex::detail::is_binary_op_supported_v>> +complex reduce_over_group(Group g, complex x, complex init, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ complex result; @@ -1770,75 +1795,91 @@ complex reduce_over_group(Group g, complex x, complex init, BinaryOpera return result; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray specialization -template > && - is_gencomplex_v && - is_gencomplex_v && - cplex::detail::is_binary_op_supported_v>> -sycl::marray reduce_over_group(Group g, sycl::marray x, sycl::marray init, BinaryOperation binary_op) { +template > && is_gencomplex_v && + is_gencomplex_v && + cplex::detail::is_binary_op_supported_v>> +sycl::marray reduce_over_group(Group g, sycl::marray x, + sycl::marray init, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; sycl::detail::loop([&](size_t s) { - result[s] = reduce_over_group(g, x[s], init[s], binary_op); + result[s] = reduce_over_group(g, x[s], init[s], binary_op); }); return result; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray and Complex specialization -template > && - (is_gencomplex_v || is_mgencomplex_v) && - cplex::detail::is_binary_op_supported_v>> +template > && + (is_gencomplex_v || is_mgencomplex_v)&&cplex::detail:: + is_binary_op_supported_v>> T reduce_over_group(Group g, T x, BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ auto init = cplex::detail::get_init(); return reduce_over_group(g, x, init, binary_op); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /* JOINT_REDUCE'S OVERLOADS */ /// Marray and Complex specialization -template > && - sycl::detail::is_pointer::value && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v || is_mgencomplex_v) && - cplex::detail::is_binary_op_supported_v>> -T joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { +template > && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v || is_mgencomplex_v)&&cplex:: + detail::is_binary_op_supported_v>> +T joint_reduce(Group g, Ptr first, Ptr last, T init, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ auto partial = cplex::detail::get_init(); - sycl::detail::for_each(g, first, last, [&](const typename sycl::detail::remove_pointer::type &x) { - partial = binary_op(partial, x); - }); + sycl::detail::for_each( + g, first, last, + [&](const typename sycl::detail::remove_pointer::type &x) { + partial = binary_op(partial, x); + }); return reduce_over_group(g, partial, init, binary_op); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray and Complex specialization -template > && - sycl::detail::is_pointer::value && - (is_gencomplex_v> || is_mgencomplex_v>) && - cplex::detail::is_binary_op_supported_v>> -typename sycl::detail::remove_pointer_t joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) { +template > && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || + is_mgencomplex_v>)&&cplex:: + detail::is_binary_op_supported_v>> +typename sycl::detail::remove_pointer_t +joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ using T = typename sycl::detail::remove_pointer_t; @@ -1846,88 +1887,109 @@ typename sycl::detail::remove_pointer_t joint_reduce(Group g, Ptr first, Pt return joint_reduce(g, first, last, init, binary_op); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /* INCLUSIVE_SCAN_OVER_GROUP'S OVERLOADS */ /// Complex specialization -template > && - is_genfloat_v && - is_genfloat_v && - cplex::detail::is_binary_op_supported_v>> -complex inclusive_scan_over_group(Group g, complex x, BinaryOperation binary_op, complex init) { +template > && is_genfloat_v && + is_genfloat_v && + cplex::detail::is_binary_op_supported_v>> +complex inclusive_scan_over_group(Group g, complex x, + BinaryOperation binary_op, + complex init) { #ifdef __SYCL_DEVICE_ONLY__ complex result; - result.real(sycl::inclusive_scan_over_group(g, x.real(), binary_op, init.real())); - result.imag(sycl::inclusive_scan_over_group(g, x.imag(), binary_op, init.imag())); + result.real( + sycl::inclusive_scan_over_group(g, x.real(), binary_op, init.real())); + result.imag( + sycl::inclusive_scan_over_group(g, x.imag(), binary_op, init.imag())); return result; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray specialization -template > && - is_gencomplex_v && - is_gencomplex_v && - cplex::detail::is_binary_op_supported_v>> -sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, BinaryOperation binary_op, sycl::marray init) { +template > && is_gencomplex_v && + is_gencomplex_v && + cplex::detail::is_binary_op_supported_v>> +sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, + BinaryOperation binary_op, + sycl::marray init) { #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; sycl::detail::loop([&](size_t s) { - result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); + result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); }); return result; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray and Complex specialization -template > && - (is_gencomplex_v || is_mgencomplex_v) && - cplex::detail::is_binary_op_supported_v>> +template > && + (is_gencomplex_v || is_mgencomplex_v)&&cplex::detail:: + is_binary_op_supported_v>> T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ auto init = cplex::detail::get_init(); return inclusive_scan_over_group(g, x, binary_op, init); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /* JOINT_INCLUSIVE_SCAN'S OVERLOADS */ /// Complex specialization -template > && - sycl::detail::is_pointer::value && - sycl::detail::is_pointer::value && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v || is_mgencomplex_v) && - cplex::detail::is_binary_op_supported_v>> -OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) { +template > && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v || + is_mgencomplex_v)&&cplex:: + detail::is_binary_op_supported_v>> +OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + BinaryOperation binary_op, T init) { #ifdef __SYCL_DEVICE_ONLY__ std::ptrdiff_t offset = g.get_local_linear_id(); std::ptrdiff_t stride = g.get_local_linear_range(); std::ptrdiff_t N = last - first; - auto roundup = [=](const std::ptrdiff_t &v, const std::ptrdiff_t &divisor) -> std::ptrdiff_t { + auto roundup = [=](const std::ptrdiff_t &v, + const std::ptrdiff_t &divisor) -> std::ptrdiff_t { return ((v + divisor - 1) / divisor) * divisor; }; - typename std::remove_const_t> x; + typename std::remove_const_t> + x; typename sycl::detail::remove_pointer_t carry = init; for (std::ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) { @@ -1936,7 +1998,8 @@ OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, Bin if (i < N) x = first[i]; - typename sycl::detail::remove_pointer_t out = inclusive_scan_over_group(g, x, binary_op, carry); + typename sycl::detail::remove_pointer_t out = + inclusive_scan_over_group(g, x, binary_op, carry); if (i < N) result[i] = out; @@ -1945,19 +2008,27 @@ OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, Bin } return result + N; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Complex specialization -template > && - sycl::detail::is_pointer::value && - sycl::detail::is_pointer::value && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v> || is_mgencomplex_v>) && - cplex::detail::is_binary_op_supported_v>> -OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { +template < + typename Group, typename InPtr, typename OutPtr, class BinaryOperation, + typename = std::enable_if_t< + sycl::is_group_v> && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v> || + is_mgencomplex_v< + sycl::detail::remove_pointer_t>)&&cplex:: + detail::is_binary_op_supported_v>> +OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ using T = typename sycl::detail::remove_pointer_t; @@ -1965,89 +2036,109 @@ OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, Bin return joint_inclusive_scan(g, first, last, result, binary_op, init); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /* EXCLUSIVE_SCAN_OVER_GROUP'S OVERLOADS */ /// Complex specialization -template > && - is_genfloat_v && - is_genfloat_v && - cplex::detail::is_binary_op_supported_v>> -complex exclusive_scan_over_group(Group g, complex x, complex init, BinaryOperation binary_op) { +template > && is_genfloat_v && + is_genfloat_v && + cplex::detail::is_binary_op_supported_v>> +complex exclusive_scan_over_group(Group g, complex x, complex init, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ complex result; - result.real(sycl::exclusive_scan_over_group(g, x.real(), init.real(), binary_op)); - result.imag(sycl::exclusive_scan_over_group(g, x.imag(), init.imag(), binary_op)); + result.real( + sycl::exclusive_scan_over_group(g, x.real(), init.real(), binary_op)); + result.imag( + sycl::exclusive_scan_over_group(g, x.imag(), init.imag(), binary_op)); return result; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray specialization -template > && - is_gencomplex_v && - is_gencomplex_v && - cplex::detail::is_binary_op_supported_v>> -sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, sycl::marray init, BinaryOperation binary_op) { +template > && is_gencomplex_v && + is_gencomplex_v && + cplex::detail::is_binary_op_supported_v>> +sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, + sycl::marray init, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; sycl::detail::loop([&](size_t s) { - result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); + result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); }); return result; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Marray and Complex specialization -template > && - (is_gencomplex_v || is_mgencomplex_v) && - cplex::detail::is_binary_op_supported_v>> +template > && + (is_gencomplex_v || is_mgencomplex_v)&&cplex::detail:: + is_binary_op_supported_v>> T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ auto init = cplex::detail::get_init(); return exclusive_scan_over_group(g, x, init, binary_op); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /* JOINT_EXCLUSIVE_SCAN'S OVERLOADS */ /// Complex specialization -template > && - sycl::detail::is_pointer::value && - sycl::detail::is_pointer::value && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v || is_mgencomplex_v) && - // - cplex::detail::is_binary_op_supported_v>> -OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) { +template > && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v || + is_mgencomplex_v)&& + // + cplex::detail::is_binary_op_supported_v>> +OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + T init, BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ std::ptrdiff_t offset = g.get_local_linear_id(); std::ptrdiff_t stride = g.get_local_linear_range(); std::ptrdiff_t N = last - first; - auto roundup = [=](const std::ptrdiff_t &v, const std::ptrdiff_t &divisor) -> std::ptrdiff_t { + auto roundup = [=](const std::ptrdiff_t &v, + const std::ptrdiff_t &divisor) -> std::ptrdiff_t { return ((v + divisor - 1) / divisor) * divisor; }; - typename std::remove_const_t> x; + typename std::remove_const_t> + x; typename sycl::detail::remove_pointer_t carry = init; for (std::ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) { @@ -2055,7 +2146,8 @@ OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T i if (i < N) x = first[i]; - typename sycl::detail::remove_pointer_t out = exclusive_scan_over_group(g, x, carry, binary_op); + typename sycl::detail::remove_pointer_t out = + exclusive_scan_over_group(g, x, carry, binary_op); if (i < N) result[i] = out; @@ -2064,19 +2156,27 @@ OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T i } return result + N; #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } /// Complex specialization -template > && - sycl::detail::is_pointer::value && - sycl::detail::is_pointer::value && - (is_gencomplex_v> || is_mgencomplex_v>) && - (is_gencomplex_v> || is_mgencomplex_v>) && - cplex::detail::is_binary_op_supported_v>> -OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { +template < + typename Group, typename InPtr, typename OutPtr, class BinaryOperation, + typename = std::enable_if_t< + sycl::is_group_v> && + sycl::detail::is_pointer::value && + sycl::detail::is_pointer::value && + (is_gencomplex_v> || + is_mgencomplex_v>)&&(is_gencomplex_v> || + is_mgencomplex_v< + sycl::detail::remove_pointer_t>)&&cplex:: + detail::is_binary_op_supported_v>> +OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, + BinaryOperation binary_op) { #ifdef __SYCL_DEVICE_ONLY__ using T = typename sycl::detail::remove_pointer_t; @@ -2084,7 +2184,8 @@ OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, Bin return joint_exclusive_scan(g, first, last, result, init, binary_op); #else - throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "Group algorithms are not supported on host."); + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Group algorithms are not supported on host."); #endif } From a1f61162771bbccde0643b8c0ce9fa130c324730 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Fri, 2 Jun 2023 11:41:10 +0200 Subject: [PATCH 09/13] format with clang-format-14 --- include/sycl_ext_complex.hpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index 22fc617..c44907d 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -305,13 +305,11 @@ template struct __numeric_type { static const bool value = !std::is_same::value; }; -template <> struct __numeric_type { - static const bool value = true; -}; +template <> struct __numeric_type { static const bool value = true; }; template ::value && __numeric_type<_A2>::value && - __numeric_type<_A3>::value> + bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value + &&__numeric_type<_A3>::value> class __promote_imp { public: static const bool value = false; From 3e83260c115077f8c1f0172224641cfa7b400d49 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Tue, 6 Jun 2023 11:29:02 +0200 Subject: [PATCH 10/13] ifdef ONEAPI for detail::loop --- include/sycl_ext_complex.hpp | 21 +++++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index c44907d..ffde5e8 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -1811,9 +1811,16 @@ sycl::marray reduce_over_group(Group g, sycl::marray x, #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; +// Use sycl::detail::loop if the compiler is intel/llvm's clang +#ifdef __SYCL_COMPILER_VERSION sycl::detail::loop([&](size_t s) { result[s] = reduce_over_group(g, x[s], init[s], binary_op); }); +#else + for (std::size_t s = 0; s < N; ++s) { + result[s] = reduce_over_group(g, x[s], init[s], binary_op); + } +#endif return result; #else @@ -1929,9 +1936,16 @@ sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; +// Use sycl::detail::loop if the compiler is intel/llvm's clang +#ifdef __SYCL_COMPILER_VERSION sycl::detail::loop([&](size_t s) { result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); }); +#else + for (std::size_t s = 0; s < N; ++s) { + result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); + } +#endif return result; #else @@ -2077,9 +2091,16 @@ sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; +// Use sycl::detail::loop if the compiler is intel/llvm's clang +#ifdef __SYCL_COMPILER_VERSION sycl::detail::loop([&](size_t s) { result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); }); +#else + for (std::size_t s = 0; s < N; ++s) { + result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); + } +#endif return result; #else From fa7fc4fcdedcbec02d85202157db1acee5eeaba2 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Tue, 6 Jun 2023 16:21:54 +0200 Subject: [PATCH 11/13] Use multiple macros to target only clang and not icpx --- include/sycl_ext_complex.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index 5989a14..91f0c27 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -1812,8 +1812,8 @@ sycl::marray reduce_over_group(Group g, sycl::marray x, sycl::marray result; // Use sycl::detail::loop if the compiler is intel/llvm's clang -#ifdef __SYCL_COMPILER_VERSION - sycl::detail::loop([&](size_t s) { +#if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) + sycl::detail::detail_loop([&](size_t s) { result[s] = reduce_over_group(g, x[s], init[s], binary_op); }); #else @@ -1937,8 +1937,8 @@ sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, sycl::marray result; // Use sycl::detail::loop if the compiler is intel/llvm's clang -#ifdef __SYCL_COMPILER_VERSION - sycl::detail::loop([&](size_t s) { +#if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) + sycl::detail::detail_loop([&](size_t s) { result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); }); #else @@ -2092,8 +2092,8 @@ sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, sycl::marray result; // Use sycl::detail::loop if the compiler is intel/llvm's clang -#ifdef __SYCL_COMPILER_VERSION - sycl::detail::loop([&](size_t s) { +#if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) + sycl::detail::detail_loop([&](size_t s) { result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); }); #else From 9fcb9feb52ba1567a51790203c27e441ee7a3d5b Mon Sep 17 00:00:00 2001 From: jle-quel Date: Tue, 6 Jun 2023 16:31:20 +0200 Subject: [PATCH 12/13] fix previous CI test - rename dim_loop to loop --- include/sycl_ext_complex.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index 91f0c27..12e1e31 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -1813,7 +1813,7 @@ sycl::marray reduce_over_group(Group g, sycl::marray x, // Use sycl::detail::loop if the compiler is intel/llvm's clang #if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) - sycl::detail::detail_loop([&](size_t s) { + sycl::detail::loop([&](size_t s) { result[s] = reduce_over_group(g, x[s], init[s], binary_op); }); #else @@ -1938,7 +1938,7 @@ sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, // Use sycl::detail::loop if the compiler is intel/llvm's clang #if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) - sycl::detail::detail_loop([&](size_t s) { + sycl::detail::loop([&](size_t s) { result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); }); #else @@ -2093,7 +2093,7 @@ sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, // Use sycl::detail::loop if the compiler is intel/llvm's clang #if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) - sycl::detail::detail_loop([&](size_t s) { + sycl::detail::loop([&](size_t s) { result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); }); #else From 41ccd28164e555c98caecfc4681adfe01c602bc8 Mon Sep 17 00:00:00 2001 From: jle-quel Date: Wed, 7 Jun 2023 16:26:23 +0200 Subject: [PATCH 13/13] implement cplex::detail's 'sycl::complex::detail::loop' --- include/sycl_ext_complex.hpp | 38 +++++++++++++----------------------- 1 file changed, 14 insertions(+), 24 deletions(-) diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp index 12e1e31..91ffde3 100644 --- a/include/sycl_ext_complex.hpp +++ b/include/sycl_ext_complex.hpp @@ -375,6 +375,17 @@ _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr bool isinf(const T a) { return sycl::isinf(a); #endif } + +// To ensure loop unrolling is done when processing dimensions. +template +void loop_impl(std::integer_sequence, F &&f) { + (f(std::integral_constant{}), ...); +} + +template void loop(F &&f) { + loop_impl(std::make_index_sequence{}, std::forward(f)); +} + } // namespace cplex::detail //////////////////////////////////////////////////////////////////////////////// @@ -1811,16 +1822,9 @@ sycl::marray reduce_over_group(Group g, sycl::marray x, #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; -// Use sycl::detail::loop if the compiler is intel/llvm's clang -#if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) - sycl::detail::loop([&](size_t s) { + cplex::detail::loop([&](size_t s) { result[s] = reduce_over_group(g, x[s], init[s], binary_op); }); -#else - for (std::size_t s = 0; s < N; ++s) { - result[s] = reduce_over_group(g, x[s], init[s], binary_op); - } -#endif return result; #else @@ -1936,16 +1940,9 @@ sycl::marray inclusive_scan_over_group(Group g, sycl::marray x, #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; -// Use sycl::detail::loop if the compiler is intel/llvm's clang -#if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) - sycl::detail::loop([&](size_t s) { + cplex::detail::loop([&](size_t s) { result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); }); -#else - for (std::size_t s = 0; s < N; ++s) { - result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]); - } -#endif return result; #else @@ -2091,16 +2088,9 @@ sycl::marray exclusive_scan_over_group(Group g, sycl::marray x, #ifdef __SYCL_DEVICE_ONLY__ sycl::marray result; -// Use sycl::detail::loop if the compiler is intel/llvm's clang -#if defined(SYCL_IMPLEMENTATION_ONEAPI) && !defined(__INTEL_LLVM_COMPILER) - sycl::detail::loop([&](size_t s) { + cplex::detail::loop([&](size_t s) { result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); }); -#else - for (std::size_t s = 0; s < N; ++s) { - result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op); - } -#endif return result; #else