From e7f792949316b4048fca6139e52e59a0dd69d629 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Fri, 31 Mar 2023 13:03:09 -0400 Subject: [PATCH 01/11] [SYCL] Add marray support to rest math built-in functions This patch adds support of sycl::marray to the rest math built-in functions (SYCL 2020, Table 175), and adds missing tests for math and common functions for https://github.com/intel/llvm/pull/8631 to reduce number of upcoming cherry-picks. --- sycl/include/sycl/builtins.hpp | 87 +++++++++++++ .../sycl/detail/generic_type_traits.hpp | 12 ++ .../DeviceLib/built-ins/marray_common.cpp | 122 ++++++++++++++++++ .../DeviceLib/built-ins/marray_math.cpp | 104 +++++++++++++++ 4 files changed, 325 insertions(+) create mode 100644 sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp create mode 100644 sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index 67083ea87d447..a1e309e4293e4 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include // TODO Decide whether to mark functions with this attribute. @@ -775,6 +776,92 @@ detail::enable_if_t::value, T> trunc(T x) __NOEXC { return __sycl_std::__invoke_trunc(x); } +// other marray math functions + +// TODO: can be optimized in the way marray math functions above are optimized +// (usage of vec) +#define __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARGPTR, \ + ...) \ + marray res; \ + for (int j = 0; j < N; j++) { \ + multi_ptr>, \ + T2::address_space> \ + ptr = nullptr; \ + res[j] = NAME(__VA_ARGS__, ptr); \ + (*ARGPTR)[j] = *ptr; \ + } \ + return res; + +#define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD( \ + NAME, ARG1, ARG2, ...) \ + template \ + std::enable_if_t::value && \ + detail::is_genfloatptr_marray::value, \ + marray> \ + NAME(marray ARG1, T2 ARG2) __NOEXC { \ + __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \ + __VA_ARGS__) \ + } + +__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(fract, x, iptr, + x[j]) +__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(modf, x, iptr, + x[j]) +__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(sincos, x, + cosval, x[j]) + +#undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENFLOATPTR_OVERLOAD + +#define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD( \ + NAME, ARG1, ARG2, ...) \ + template \ + std::enable_if_t::value && \ + detail::is_genintptr_marray::value, \ + marray> \ + NAME(marray ARG1, T2 ARG2) __NOEXC { \ + __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \ + __VA_ARGS__) \ + } + +__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(frexp, x, exp, + x[j]) +__SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(lgamma_r, x, signp, + x[j]) + +#undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENINTPTR_OVERLOAD + +#define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME, ...) \ + template \ + std::enable_if_t::value && \ + detail::is_genintptr_marray::value, \ + marray> \ + NAME(marray x, marray y, T2 quo) __NOEXC { \ + __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, quo, \ + __VA_ARGS__) \ + } + +__SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(remquo, x[j], y[j]) + +#undef __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD + +#undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL + +#define __SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD(NAME) \ + template \ + std::enable_if_t::value, \ + marray, N>> \ + NAME(marray nancode) __NOEXC { \ + marray res; \ + for (int j = 0; j < N; j++) { \ + res[j] = NAME(nancode[j]); \ + } \ + return res; \ + } + +__SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD(nan) + +#undef __SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD + /* --------------- 4.13.5 Common functions. ---------------------------------*/ // svgenfloat clamp (svgenfloat x, svgenfloat minval, svgenfloat maxval) template diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp index 62fe70e0491e9..f6dd7459e8544 100644 --- a/sycl/include/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/sycl/detail/generic_type_traits.hpp @@ -232,11 +232,23 @@ using is_genintptr = bool_constant< is_pointer::value && is_genint>::value && is_address_space_compliant::value>; +template +using is_genintptr_marray = bool_constant< + is_pointer::value && + is_genint>>::value && + is_address_space_compliant::value>; + template using is_genfloatptr = bool_constant< is_pointer::value && is_genfloat>::value && is_address_space_compliant::value>; +template +using is_genfloatptr_marray = bool_constant< + is_pointer::value && + is_genfloat>>::value && + is_address_space_compliant::value>; + template using is_genptr = bool_constant< is_pointer::value && is_gentype>::value && diff --git a/sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp b/sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp new file mode 100644 index 0000000000000..76e02a9f53e40 --- /dev/null +++ b/sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp @@ -0,0 +1,122 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +#ifdef _WIN32 +#define _USE_MATH_DEFINES // To use math constants +#include +#endif + +#include + +#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \ + { \ + { \ + MARRAY_ELEM_TYPE result[DIM]; \ + { \ + sycl::buffer b(result, sycl::range{DIM}); \ + deviceQueue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { \ + sycl::marray res = FUNC(__VA_ARGS__); \ + for (int i = 0; i < DIM; i++) \ + res_access[i] = res[i]; \ + }); \ + }); \ + } \ + for (int i = 0; i < DIM; i++) \ + assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ + } \ + } + +#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) + +int main() { + sycl::queue deviceQueue; + sycl::device dev = deviceQueue.get_device(); + + sycl::marray ma1{1.0f, 2.0f}; + sycl::marray ma2{1.0f, 2.0f}; + sycl::marray ma3{3.0f, 2.0f}; + sycl::marray ma4{1.0, 2.0}; + sycl::marray ma5{M_PI, M_PI, M_PI}; + sycl::marray ma6{M_PI, M_PI, M_PI}; + sycl::marray ma7{M_PI, M_PI, M_PI}; + sycl::marray ma8{0.3f, 0.6f}; + sycl::marray ma9{5.0, 8.0}; + sycl::marray ma10{180, 180, 180}; + sycl::marray ma11{180, 180, 180}; + sycl::marray ma12{180, 180, 180}; + sycl::marray ma13{181, 179, 181}; + sycl::marray ma14{+0.0f, -0.6f}; + sycl::marray ma15{-0.0, 0.6f}; + + // sycl::clamp + TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2, ma3); + TEST(sycl::clamp, float, 2, EXPECTED(float, 2.0f, 2.0f), 0, ma1, 3.0f, 2.0f); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::clamp, double, 2, EXPECTED(double, 2.0, 2.0), 0, ma4, 3.0, 2.0); + // sycl::degrees + TEST(sycl::degrees, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::degrees, double, 3, EXPECTED(double, 180, 180, 180), 0, ma6); + if (dev.has(sycl::aspect::fp16)) + TEST(sycl::degrees, sycl::half, 3, EXPECTED(sycl::half, 180, 180, 180), 0.2, + ma7); + // sycl::max + TEST(sycl::max, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma3); + TEST(sycl::max, float, 2, EXPECTED(float, 1.5f, 2.0f), 0, ma1, 1.5f); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::max, double, 2, EXPECTED(double, 1.5, 2.0), 0, ma4, 1.5); + // sycl::min + TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma3); + TEST(sycl::min, float, 2, EXPECTED(float, 1.0f, 1.5f), 0, ma1, 1.5f); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::min, double, 2, EXPECTED(double, 1.0, 1.5), 0, ma4, 1.5); + // sycl::mix + TEST(sycl::mix, float, 2, EXPECTED(float, 1.6f, 2.0f), 0, ma1, ma3, ma8); + TEST(sycl::mix, float, 2, EXPECTED(float, 1.4f, 2.0f), 0, ma1, ma3, 0.2); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::mix, double, 2, EXPECTED(double, 3.0, 5.0), 0, ma4, ma9, 0.5); + // sycl::radians + TEST(sycl::radians, float, 3, EXPECTED(float, M_PI, M_PI, M_PI), 0, ma10); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::radians, double, 3, EXPECTED(double, M_PI, M_PI, M_PI), 0, ma11); + if (dev.has(sycl::aspect::fp16)) + TEST(sycl::radians, sycl::half, 3, EXPECTED(sycl::half, M_PI, M_PI, M_PI), + 0.002, ma12); + // sycl::step + TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 1.0f), 0, ma1, ma3); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::step, double, 2, EXPECTED(double, 1.0, 1.0), 0, ma4, ma9); + if (dev.has(sycl::aspect::fp16)) + TEST(sycl::step, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0), 0, + ma12, ma13); + TEST(sycl::step, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, 2.5f, ma3); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::step, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, ma9); + // sycl::smoothstep + TEST(sycl::smoothstep, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, ma1, ma2, + ma3); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.00147576, 0.00446826), + 0.00000001, ma4, ma11, ma9); + if (dev.has(sycl::aspect::fp16)) + TEST(sycl::smoothstep, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0), + 0, ma12, ma12, ma13); + TEST(sycl::smoothstep, float, 2, EXPECTED(float, 0.0553936f, 0.0f), 0.0000001, + 2.5f, 6.0f, ma3); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, + 8.0f, ma9); + // sign + TEST(sycl::sign, float, 2, EXPECTED(float, +0.0f, -1.0f), 0, ma14); + if (dev.has(sycl::aspect::fp64)) + TEST(sycl::sign, double, 2, EXPECTED(double, -0.0, 1.0), 0, ma15); + if (dev.has(sycl::aspect::fp16)) + TEST(sycl::sign, sycl::half, 3, EXPECTED(sycl::half, 1.0, 1.0, 1.0), 0, + ma12); + + return 0; +} diff --git a/sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp b/sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp new file mode 100644 index 0000000000000..8d38d60ea04e4 --- /dev/null +++ b/sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp @@ -0,0 +1,104 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +#include + +#define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...) \ + { \ + { \ + MARRAY_ELEM_TYPE result[DIM]; \ + { \ + sycl::buffer b(result, sycl::range{DIM}); \ + deviceQueue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { \ + sycl::marray res = FUNC(__VA_ARGS__); \ + for (int i = 0; i < DIM; i++) \ + res_access[i] = res[i]; \ + }); \ + }); \ + } \ + for (int i = 0; i < DIM; i++) \ + assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ + } \ + } + +#define TEST2(FUNC, MARRAY_ELEM_TYPE, PTR_TYPE, DIM, EXPECTED_1, EXPECTED_2, \ + DELTA, ...) \ + { \ + { \ + MARRAY_ELEM_TYPE result[DIM]; \ + sycl::marray result_ptr; \ + { \ + sycl::buffer b(result, sycl::range{DIM}); \ + sycl::buffer, 1> b_ptr(&result_ptr, \ + sycl::range<1>(1)); \ + deviceQueue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + sycl::accessor res_ptr_access{b_ptr, cgh}; \ + cgh.single_task([=]() { \ + sycl::global_ptr> ptr(res_ptr_access); \ + sycl::marray res = FUNC(__VA_ARGS__, ptr); \ + for (int i = 0; i < DIM; i++) \ + res_access[i] = res[i]; \ + }); \ + }); \ + } \ + for (int i = 0; i < DIM; i++) { \ + assert(std::abs(result[i] - EXPECTED_1[i]) <= DELTA); \ + assert(std::abs(result_ptr[i] - EXPECTED_2[i]) <= DELTA); \ + } \ + } \ + } + +#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) + +int main() { + sycl::queue deviceQueue; + + sycl::marray ma1{1.0f, 2.0f}; + sycl::marray ma2{3.0f, 2.0f}; + sycl::marray ma3{180, 180, 180}; + sycl::marray ma4{1, 1, 1}; + sycl::marray ma5{180, -180, -180}; + sycl::marray ma6{1.4f, 4.2f, 5.3f}; + sycl::marray ma7{1, 2, 3}; + sycl::marray ma8{1, 2, 3}; + + TEST(sycl::fabs, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5); + TEST(sycl::ilogb, int, 3, EXPECTED(int, 7, 7, 7), 0, ma3); + TEST(sycl::fmax, float, 2, EXPECTED(float, 3.0f, 2.0f), 0, ma1, ma2); + TEST(sycl::fmax, float, 2, EXPECTED(float, 5.0f, 5.0f), 0, ma1, 5.0f); + TEST(sycl::fmin, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2); + TEST(sycl::fmin, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, 5.0f); + TEST(sycl::ldexp, float, 3, EXPECTED(float, 360, 360, 360), 0, ma3, ma4); + TEST(sycl::ldexp, float, 3, EXPECTED(float, 5760, 5760, 5760), 0, ma3, 5); + TEST(sycl::pown, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, ma4); + TEST(sycl::pown, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, 1); + TEST(sycl::rootn, float, 3, EXPECTED(float, 180, 180, 180), 0.1, ma3, ma4); + TEST(sycl::rootn, float, 3, EXPECTED(float, 2.82523, 2.82523, 2.82523), + 0.00001, ma3, 5); + TEST2(sycl::fract, float, float, 3, EXPECTED(float, 0.4f, 0.2f, 0.3f), + EXPECTED(float, 1, 4, 5), 0.0001, ma6); + TEST2(sycl::modf, float, float, 3, EXPECTED(float, 0.4f, 0.2f, 0.3f), + EXPECTED(float, 1, 4, 5), 0.0001, ma6); + TEST2(sycl::sincos, float, float, 3, + EXPECTED(float, 0.98545f, -0.871576f, -0.832267f), + EXPECTED(float, 0.169967, -0.490261, 0.554375), 0.0001, ma6); + TEST2(sycl::frexp, float, int, 3, EXPECTED(float, 0.7f, 0.525f, 0.6625f), + EXPECTED(int, 1, 3, 3), 0.0001, ma6); + TEST2(sycl::lgamma_r, float, int, 3, + EXPECTED(float, -0.119613f, 2.04856f, 3.63964f), EXPECTED(int, 1, 1, 1), + 0.0001, ma6); + TEST2(sycl::remquo, float, int, 3, EXPECTED(float, 1.4f, 4.2f, 5.3f), + EXPECTED(int, 0, 0, 0), 0.0001, ma6, ma3); + TEST(sycl::nan, float, 3, EXPECTED(float, 0, 0, 0), 0.1, ma7); + TEST(sycl::nan, double, 3, EXPECTED(double, 0, 0, 0), 0.1, ma8); + + TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1, + ma1); + + return 0; +} From 100a872691bf000d93128f138bbc48ff8d73ba50 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Tue, 4 Apr 2023 07:22:19 -0400 Subject: [PATCH 02/11] Apply CR comment, fixed other bugs --- sycl/include/sycl/builtins.hpp | 71 ++++++++++--------- .../sycl/detail/generic_type_traits.hpp | 21 ++++-- .../built-ins/marray_common.cpp | 4 +- .../built-ins/marray_math.cpp | 41 ++++++++++- 4 files changed, 91 insertions(+), 46 deletions(-) rename sycl/test-e2e/{DeviceLib => Basic}/built-ins/marray_common.cpp (97%) rename sycl/test-e2e/{DeviceLib => Basic}/built-ins/marray_math.cpp (69%) diff --git a/sycl/include/sycl/builtins.hpp b/sycl/include/sycl/builtins.hpp index a1e309e4293e4..0519736d8d59e 100644 --- a/sycl/include/sycl/builtins.hpp +++ b/sycl/include/sycl/builtins.hpp @@ -784,21 +784,23 @@ detail::enable_if_t::value, T> trunc(T x) __NOEXC { ...) \ marray res; \ for (int j = 0; j < N; j++) { \ - multi_ptr>, \ - T2::address_space> \ - ptr = nullptr; \ - res[j] = NAME(__VA_ARGS__, ptr); \ - (*ARGPTR)[j] = *ptr; \ + res[j] = \ + NAME(__VA_ARGS__, \ + address_space_cast>(&(*ARGPTR)[j])); \ } \ return res; #define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD( \ NAME, ARG1, ARG2, ...) \ - template \ - std::enable_if_t::value && \ - detail::is_genfloatptr_marray::value, \ - marray> \ - NAME(marray ARG1, T2 ARG2) __NOEXC { \ + template \ + std::enable_if_t< \ + detail::is_svgenfloat::value && \ + detail::is_genfloatptr_marray::value, \ + marray> \ + NAME(marray ARG1, multi_ptr ARG2) \ + __NOEXC { \ __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \ __VA_ARGS__) \ } @@ -814,11 +816,14 @@ __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENFLOATPTR_OVERLOAD(sincos, x, #define __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD( \ NAME, ARG1, ARG2, ...) \ - template \ - std::enable_if_t::value && \ - detail::is_genintptr_marray::value, \ - marray> \ - NAME(marray ARG1, T2 ARG2) __NOEXC { \ + template \ + std::enable_if_t< \ + detail::is_svgenfloat::value && \ + detail::is_genintptr_marray::value, \ + marray> \ + NAME(marray ARG1, multi_ptr ARG2) \ + __NOEXC { \ __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, ARG2, \ __VA_ARGS__) \ } @@ -831,11 +836,14 @@ __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_ARG_GENINTPTR_OVERLOAD(lgamma_r, x, signp, #undef __SYCL_MARRAY_MATH_FUNCTION_BINOP_2ND_GENINTPTR_OVERLOAD #define __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(NAME, ...) \ - template \ - std::enable_if_t::value && \ - detail::is_genintptr_marray::value, \ - marray> \ - NAME(marray x, marray y, T2 quo) __NOEXC { \ + template \ + std::enable_if_t< \ + detail::is_svgenfloat::value && \ + detail::is_genintptr_marray::value, \ + marray> \ + NAME(marray x, marray y, \ + multi_ptr quo) __NOEXC { \ __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL(NAME, quo, \ __VA_ARGS__) \ } @@ -846,21 +854,16 @@ __SYCL_MARRAY_MATH_FUNCTION_REMQUO_OVERLOAD(remquo, x[j], y[j]) #undef __SYCL_MARRAY_MATH_FUNCTION_W_GENPTR_ARG_OVERLOAD_IMPL -#define __SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD(NAME) \ - template \ - std::enable_if_t::value, \ - marray, N>> \ - NAME(marray nancode) __NOEXC { \ - marray res; \ - for (int j = 0; j < N; j++) { \ - res[j] = NAME(nancode[j]); \ - } \ - return res; \ +template +std::enable_if_t::value, + marray, N>> +nan(marray nancode) __NOEXC { + marray, N> res; + for (int j = 0; j < N; j++) { + res[j] = nan(nancode[j]); } - -__SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD(nan) - -#undef __SYCL_MARRAY_MATH_FUNCTION_NAN_OVERLOAD + return res; +} /* --------------- 4.13.5 Common functions. ---------------------------------*/ // svgenfloat clamp (svgenfloat x, svgenfloat minval, svgenfloat maxval) diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp index f6dd7459e8544..a04ce83092c86 100644 --- a/sycl/include/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/sycl/detail/generic_type_traits.hpp @@ -232,22 +232,29 @@ using is_genintptr = bool_constant< is_pointer::value && is_genint>::value && is_address_space_compliant::value>; -template +template using is_genintptr_marray = bool_constant< - is_pointer::value && + std::is_same, T::size()>>::value && is_genint>>::value && - is_address_space_compliant::value>; + is_address_space_compliant, + gvl::nonconst_address_space_list>::value && + (IsDecorated == access::decorated::yes || + IsDecorated == access::decorated::no)>; template using is_genfloatptr = bool_constant< is_pointer::value && is_genfloat>::value && is_address_space_compliant::value>; -template +template using is_genfloatptr_marray = bool_constant< - is_pointer::value && - is_genfloat>>::value && - is_address_space_compliant::value>; + is_mgenfloat::value && + is_address_space_compliant, + gvl::nonconst_address_space_list>::value && + (IsDecorated == access::decorated::yes || + IsDecorated == access::decorated::no)>; template using is_genptr = bool_constant< diff --git a/sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp b/sycl/test-e2e/Basic/built-ins/marray_common.cpp similarity index 97% rename from sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp rename to sycl/test-e2e/Basic/built-ins/marray_common.cpp index 76e02a9f53e40..0bbe20c070e5a 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/marray_common.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_common.cpp @@ -54,9 +54,9 @@ int main() { // sycl::clamp TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2, ma3); - TEST(sycl::clamp, float, 2, EXPECTED(float, 2.0f, 2.0f), 0, ma1, 3.0f, 2.0f); + TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, 1.0f, 3.0f); if (dev.has(sycl::aspect::fp64)) - TEST(sycl::clamp, double, 2, EXPECTED(double, 2.0, 2.0), 0, ma4, 3.0, 2.0); + TEST(sycl::clamp, double, 2, EXPECTED(double, 1.0, 2.0), 0, ma4, 1.0, 3.0); // sycl::degrees TEST(sycl::degrees, float, 3, EXPECTED(float, 180, 180, 180), 0, ma5); if (dev.has(sycl::aspect::fp64)) diff --git a/sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp similarity index 69% rename from sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp rename to sycl/test-e2e/Basic/built-ins/marray_math.cpp index 8d38d60ea04e4..8fe1c9c45c534 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -39,7 +39,10 @@ sycl::accessor res_access{b, cgh}; \ sycl::accessor res_ptr_access{b_ptr, cgh}; \ cgh.single_task([=]() { \ - sycl::global_ptr> ptr(res_ptr_access); \ + sycl::multi_ptr, \ + sycl::access::address_space::global_space, \ + sycl::access::decorated::yes> \ + ptr(res_ptr_access); \ sycl::marray res = FUNC(__VA_ARGS__, ptr); \ for (int i = 0; i < DIM; i++) \ res_access[i] = res[i]; \ @@ -53,6 +56,33 @@ } \ } +#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...) \ + { \ + { \ + MARRAY_ELEM_TYPE result[DIM]; \ + { \ + sycl::buffer b(result, sycl::range{DIM}); \ + deviceQueue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { \ + sycl::marray res = FUNC(__VA_ARGS__); \ + for (int i = 0; i < DIM; i++) \ + res_access[i] = res[i]; \ + }); \ + }); \ + } \ + for (int i = 0; i < DIM; i++) { \ + std::uint64_t result_uint64; \ + std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ + std::ostringstream stream; \ + stream << "0x" << std::hex << result_uint64; \ + std::string result_string = stream.str(); \ + result_string = result_string.substr(result_string.size() - 5); \ + assert(result_string.compare(EXPECTED[i]) == 0); \ + } \ + } \ + } + #define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) int main() { @@ -94,8 +124,13 @@ int main() { 0.0001, ma6); TEST2(sycl::remquo, float, int, 3, EXPECTED(float, 1.4f, 4.2f, 5.3f), EXPECTED(int, 0, 0, 0), 0.0001, ma6, ma3); - TEST(sycl::nan, float, 3, EXPECTED(float, 0, 0, 0), 0.1, ma7); - TEST(sycl::nan, double, 3, EXPECTED(double, 0, 0, 0), 0.1, ma8); + + if (!deviceQueue.get_device().is_gpu()) { + TEST3(sycl::nan, float, 3, EXPECTED(std::string, "00001", "00002", "00003"), + ma7); + TEST3(sycl::nan, double, 3, + EXPECTED(std::string, "00001", "00002", "00003"), ma8); + } TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1, ma1); From 59fb71a935315474917b237ec63d10891d81ab3b Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Tue, 4 Apr 2023 08:41:23 -0400 Subject: [PATCH 03/11] Debugging test failures --- sycl/test-e2e/Basic/built-ins/marray_common.cpp | 2 +- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 8 ++++---- 2 files changed, 5 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_common.cpp b/sycl/test-e2e/Basic/built-ins/marray_common.cpp index 0bbe20c070e5a..56c60ed80c3da 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_common.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_common.cpp @@ -26,7 +26,7 @@ }); \ } \ for (int i = 0; i < DIM; i++) \ - assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ + std::cout << result[i] << " " << EXPECTED[i] << std::endl; \ } \ } diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index 8fe1c9c45c534..0f536cd59968d 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -21,7 +21,7 @@ }); \ } \ for (int i = 0; i < DIM; i++) \ - assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ + std::cout << result[i] << " " << EXPECTED[i] << std::endl; \ } \ } @@ -50,8 +50,8 @@ }); \ } \ for (int i = 0; i < DIM; i++) { \ - assert(std::abs(result[i] - EXPECTED_1[i]) <= DELTA); \ - assert(std::abs(result_ptr[i] - EXPECTED_2[i]) <= DELTA); \ + std::cout << result[i] << " " << EXPECTED_1[i] << std::endl; \ + std::cout << result_ptr[i] << " " << EXPECTED_2[i] << std::endl; \ } \ } \ } @@ -78,7 +78,7 @@ stream << "0x" << std::hex << result_uint64; \ std::string result_string = stream.str(); \ result_string = result_string.substr(result_string.size() - 5); \ - assert(result_string.compare(EXPECTED[i]) == 0); \ + std::cout << result_string << " " << EXPECTED[i] << std::endl; \ } \ } \ } From 8bac796d78dcdbc408f4d7ef9c2085fc121bc79f Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Tue, 4 Apr 2023 09:48:57 -0400 Subject: [PATCH 04/11] Debugging test failures --- .../Basic/built-ins/marray_common.cpp | 2 + sycl/test-e2e/Basic/built-ins/marray_math.cpp | 59 ++++++++++++------- 2 files changed, 39 insertions(+), 22 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_common.cpp b/sycl/test-e2e/Basic/built-ins/marray_common.cpp index 56c60ed80c3da..e9ac6054668c3 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_common.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_common.cpp @@ -27,6 +27,8 @@ } \ for (int i = 0; i < DIM; i++) \ std::cout << result[i] << " " << EXPECTED[i] << std::endl; \ + for (int i = 0; i < DIM; i++) \ + assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ } \ } diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index 0f536cd59968d..1324b0d625003 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -22,6 +22,8 @@ } \ for (int i = 0; i < DIM; i++) \ std::cout << result[i] << " " << EXPECTED[i] << std::endl; \ + for (int i = 0; i < DIM; i++) \ + assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ } \ } @@ -53,35 +55,48 @@ std::cout << result[i] << " " << EXPECTED_1[i] << std::endl; \ std::cout << result_ptr[i] << " " << EXPECTED_2[i] << std::endl; \ } \ + for (int i = 0; i < DIM; i++) { \ + assert(std::abs(result[i] - EXPECTED_1[i]) <= DELTA); \ + assert(std::abs(result_ptr[i] - EXPECTED_2[i]) <= DELTA); \ + } } \ } -#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...) \ - { \ +#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...) \ { \ - MARRAY_ELEM_TYPE result[DIM]; \ { \ - sycl::buffer b(result, sycl::range{DIM}); \ - deviceQueue.submit([&](sycl::handler &cgh) { \ - sycl::accessor res_access{b, cgh}; \ - cgh.single_task([=]() { \ - sycl::marray res = FUNC(__VA_ARGS__); \ - for (int i = 0; i < DIM; i++) \ - res_access[i] = res[i]; \ + MARRAY_ELEM_TYPE result[DIM]; \ + { \ + sycl::buffer b(result, sycl::range{DIM}); \ + deviceQueue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { \ + sycl::marray res = FUNC(__VA_ARGS__); \ + for (int i = 0; i < DIM; i++) \ + res_access[i] = res[i]; \ + }); \ }); \ - }); \ + } \ + for (int i = 0; i < DIM; i++) { \ + std::uint64_t result_uint64; \ + std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ + std::ostringstream stream; \ + stream << "0x" << std::hex << result_uint64; \ + std::string result_string = stream.str(); \ + result_string = result_string.substr(result_string.size() - 5); \ + std::cout << result_string << " " << EXPECTED[i] << std::endl; \ + } \ + for (int i = 0; i < DIM; i++) { \ + std::uint64_t result_uint64; \ + std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ + std::ostringstream stream; \ + stream << "0x" << std::hex << result_uint64; \ + std::string result_string = stream.str(); \ + result_string = result_string.substr(result_string.size() - 5); \ + assert(result_string.compare(EXPECTED[i]) == 0); \ + } \ } \ - for (int i = 0; i < DIM; i++) { \ - std::uint64_t result_uint64; \ - std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ - std::ostringstream stream; \ - stream << "0x" << std::hex << result_uint64; \ - std::string result_string = stream.str(); \ - result_string = result_string.substr(result_string.size() - 5); \ - std::cout << result_string << " " << EXPECTED[i] << std::endl; \ - } \ - } \ - } + } #define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) From f8c049dc7d3151763606d2150890f438d88e7b91 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Tue, 4 Apr 2023 11:16:21 -0400 Subject: [PATCH 05/11] Debugging test failures --- .../Basic/built-ins/marray_common.cpp | 12 ++-- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 64 +++++++++---------- 2 files changed, 38 insertions(+), 38 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_common.cpp b/sycl/test-e2e/Basic/built-ins/marray_common.cpp index e9ac6054668c3..e64be62e55ff7 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_common.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_common.cpp @@ -99,14 +99,14 @@ int main() { if (dev.has(sycl::aspect::fp64)) TEST(sycl::step, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, ma9); // sycl::smoothstep - TEST(sycl::smoothstep, float, 2, EXPECTED(float, 1.0f, 0.0f), 0, ma1, ma2, - ma3); + TEST(sycl::smoothstep, float, 2, EXPECTED(float, 1.0f, 1.0f), 0, ma8, ma1, + ma2); if (dev.has(sycl::aspect::fp64)) - TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.00147576, 0.00446826), - 0.00000001, ma4, ma11, ma9); + TEST(sycl::smoothstep, double, 2, EXPECTED(double, 1.0, 1.0f), 0.00000001, + ma4, ma9, ma9); if (dev.has(sycl::aspect::fp16)) - TEST(sycl::smoothstep, sycl::half, 3, EXPECTED(sycl::half, 1.0, 0.0, 1.0), - 0, ma12, ma12, ma13); + TEST(sycl::smoothstep, sycl::half, 3, EXPECTED(sycl::half, 1.0, 1.0, 1.0), + 0, ma7, ma12, ma13); TEST(sycl::smoothstep, float, 2, EXPECTED(float, 0.0553936f, 0.0f), 0.0000001, 2.5f, 6.0f, ma3); if (dev.has(sycl::aspect::fp64)) diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index 1324b0d625003..821ddde3f4d8f 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -58,45 +58,45 @@ for (int i = 0; i < DIM; i++) { \ assert(std::abs(result[i] - EXPECTED_1[i]) <= DELTA); \ assert(std::abs(result_ptr[i] - EXPECTED_2[i]) <= DELTA); \ - } + } \ } \ } -#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...) \ +#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...) \ + { \ { \ + MARRAY_ELEM_TYPE result[DIM]; \ { \ - MARRAY_ELEM_TYPE result[DIM]; \ - { \ - sycl::buffer b(result, sycl::range{DIM}); \ - deviceQueue.submit([&](sycl::handler &cgh) { \ - sycl::accessor res_access{b, cgh}; \ - cgh.single_task([=]() { \ - sycl::marray res = FUNC(__VA_ARGS__); \ - for (int i = 0; i < DIM; i++) \ - res_access[i] = res[i]; \ - }); \ + sycl::buffer b(result, sycl::range{DIM}); \ + deviceQueue.submit([&](sycl::handler &cgh) { \ + sycl::accessor res_access{b, cgh}; \ + cgh.single_task([=]() { \ + sycl::marray res = FUNC(__VA_ARGS__); \ + for (int i = 0; i < DIM; i++) \ + res_access[i] = res[i]; \ }); \ - } \ - for (int i = 0; i < DIM; i++) { \ - std::uint64_t result_uint64; \ - std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ - std::ostringstream stream; \ - stream << "0x" << std::hex << result_uint64; \ - std::string result_string = stream.str(); \ - result_string = result_string.substr(result_string.size() - 5); \ - std::cout << result_string << " " << EXPECTED[i] << std::endl; \ - } \ - for (int i = 0; i < DIM; i++) { \ - std::uint64_t result_uint64; \ - std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ - std::ostringstream stream; \ - stream << "0x" << std::hex << result_uint64; \ - std::string result_string = stream.str(); \ - result_string = result_string.substr(result_string.size() - 5); \ - assert(result_string.compare(EXPECTED[i]) == 0); \ - } \ + }); \ + } \ + for (int i = 0; i < DIM; i++) { \ + std::uint64_t result_uint64; \ + std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ + std::ostringstream stream; \ + stream << "0x" << std::hex << result_uint64; \ + std::string result_string = stream.str(); \ + result_string = result_string.substr(result_string.size() - 5); \ + std::cout << result_string << " " << EXPECTED[i] << std::endl; \ } \ - } + for (int i = 0; i < DIM; i++) { \ + std::uint64_t result_uint64; \ + std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ + std::ostringstream stream; \ + stream << "0x" << std::hex << result_uint64; \ + std::string result_string = stream.str(); \ + result_string = result_string.substr(result_string.size() - 5); \ + assert(result_string.compare(EXPECTED[i]) == 0); \ + } \ + } \ + } #define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) From d745d4b9196eb39f1ceac34d6f3b4e514fe8e0bc Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Tue, 4 Apr 2023 23:19:51 +0200 Subject: [PATCH 06/11] Debugging test failures --- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index 821ddde3f4d8f..b6b5033a7c336 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -85,14 +85,6 @@ std::string result_string = stream.str(); \ result_string = result_string.substr(result_string.size() - 5); \ std::cout << result_string << " " << EXPECTED[i] << std::endl; \ - } \ - for (int i = 0; i < DIM; i++) { \ - std::uint64_t result_uint64; \ - std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ - std::ostringstream stream; \ - stream << "0x" << std::hex << result_uint64; \ - std::string result_string = stream.str(); \ - result_string = result_string.substr(result_string.size() - 5); \ assert(result_string.compare(EXPECTED[i]) == 0); \ } \ } \ From ae65305210b819812752fd7db1acad4a467e20de Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Tue, 4 Apr 2023 23:21:04 +0200 Subject: [PATCH 07/11] Debugging test failures --- sycl/test-e2e/Basic/built-ins/marray_common.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_common.cpp b/sycl/test-e2e/Basic/built-ins/marray_common.cpp index e64be62e55ff7..91fcc1ae9ba00 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_common.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_common.cpp @@ -25,8 +25,6 @@ }); \ }); \ } \ - for (int i = 0; i < DIM; i++) \ - std::cout << result[i] << " " << EXPECTED[i] << std::endl; \ for (int i = 0; i < DIM; i++) \ assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ } \ From f58d6842b715176f6704208fde2ce0bff32e0a46 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Wed, 5 Apr 2023 21:31:02 +0200 Subject: [PATCH 08/11] Fix marray math builtins API test --- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 19 ++++--------------- 1 file changed, 4 insertions(+), 15 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index b6b5033a7c336..3aa6061572c86 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -20,8 +20,6 @@ }); \ }); \ } \ - for (int i = 0; i < DIM; i++) \ - std::cout << result[i] << " " << EXPECTED[i] << std::endl; \ for (int i = 0; i < DIM; i++) \ assert(abs(result[i] - EXPECTED[i]) <= DELTA); \ } \ @@ -51,10 +49,6 @@ }); \ }); \ } \ - for (int i = 0; i < DIM; i++) { \ - std::cout << result[i] << " " << EXPECTED_1[i] << std::endl; \ - std::cout << result_ptr[i] << " " << EXPECTED_2[i] << std::endl; \ - } \ for (int i = 0; i < DIM; i++) { \ assert(std::abs(result[i] - EXPECTED_1[i]) <= DELTA); \ assert(std::abs(result_ptr[i] - EXPECTED_2[i]) <= DELTA); \ @@ -84,7 +78,6 @@ stream << "0x" << std::hex << result_uint64; \ std::string result_string = stream.str(); \ result_string = result_string.substr(result_string.size() - 5); \ - std::cout << result_string << " " << EXPECTED[i] << std::endl; \ assert(result_string.compare(EXPECTED[i]) == 0); \ } \ } \ @@ -131,14 +124,10 @@ int main() { 0.0001, ma6); TEST2(sycl::remquo, float, int, 3, EXPECTED(float, 1.4f, 4.2f, 5.3f), EXPECTED(int, 0, 0, 0), 0.0001, ma6, ma3); - - if (!deviceQueue.get_device().is_gpu()) { - TEST3(sycl::nan, float, 3, EXPECTED(std::string, "00001", "00002", "00003"), - ma7); - TEST3(sycl::nan, double, 3, - EXPECTED(std::string, "00001", "00002", "00003"), ma8); - } - + TEST3(sycl::nan, float, 3, EXPECTED(std::string, "fffff", "fffff", "fffff"), + ma7); + TEST3(sycl::nan, double, 3, EXPECTED(std::string, "fffff", "fffff", "fffff"), + ma8); TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1, ma1); From 6acfc04e1866320a079e3d68792244ec16a583f1 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Wed, 5 Apr 2023 23:27:38 +0200 Subject: [PATCH 09/11] Update marray_math.cpp Debugging test failures --- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index 3aa6061572c86..c9cdd4f7be862 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -78,6 +78,7 @@ stream << "0x" << std::hex << result_uint64; \ std::string result_string = stream.str(); \ result_string = result_string.substr(result_string.size() - 5); \ + std::cout << result_string << " " << EXPECTED[i] << std::endl; \ assert(result_string.compare(EXPECTED[i]) == 0); \ } \ } \ From 052e0cccf191335edf838142b2cb73c5bd142e6f Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Thu, 6 Apr 2023 12:59:40 +0200 Subject: [PATCH 10/11] Fix marray_math.cpp --- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index c9cdd4f7be862..8c0ae002d5874 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -125,10 +125,16 @@ int main() { 0.0001, ma6); TEST2(sycl::remquo, float, int, 3, EXPECTED(float, 1.4f, 4.2f, 5.3f), EXPECTED(int, 0, 0, 0), 0.0001, ma6, ma3); - TEST3(sycl::nan, float, 3, EXPECTED(std::string, "fffff", "fffff", "fffff"), - ma7); - TEST3(sycl::nan, double, 3, EXPECTED(std::string, "fffff", "fffff", "fffff"), - ma8); + auto backend = deviceQueue.get_device().get_backend(); + // TODO: enable for all backends when OpenCL CPU and OpenCL/Level Zero GPU + // return correct results for nan function + if (backend == sycl::backend::ext_oneapi_cuda || + backend == sycl::backend::ext_oneapi_hip) { + TEST3(sycl::nan, float, 3, EXPECTED(std::string, "00001", "00002", "00003"), + ma7); + TEST3(sycl::nan, double, 3, + EXPECTED(std::string, "00001", "00002", "00003"), ma8); + } TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1, ma1); From 2814d7ef9ffb65755a29f8fcddc437b9a9aadf85 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Thu, 6 Apr 2023 14:51:22 +0200 Subject: [PATCH 11/11] Update marray_math.cpp --- sycl/test-e2e/Basic/built-ins/marray_math.cpp | 24 +++---------------- 1 file changed, 3 insertions(+), 21 deletions(-) diff --git a/sycl/test-e2e/Basic/built-ins/marray_math.cpp b/sycl/test-e2e/Basic/built-ins/marray_math.cpp index 8c0ae002d5874..8f4a69822003f 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_math.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_math.cpp @@ -56,7 +56,7 @@ } \ } -#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...) \ +#define TEST3(FUNC, MARRAY_ELEM_TYPE, DIM, ...) \ { \ { \ MARRAY_ELEM_TYPE result[DIM]; \ @@ -71,16 +71,6 @@ }); \ }); \ } \ - for (int i = 0; i < DIM; i++) { \ - std::uint64_t result_uint64; \ - std::memcpy(&result_uint64, &result[i], sizeof(result[i])); \ - std::ostringstream stream; \ - stream << "0x" << std::hex << result_uint64; \ - std::string result_string = stream.str(); \ - result_string = result_string.substr(result_string.size() - 5); \ - std::cout << result_string << " " << EXPECTED[i] << std::endl; \ - assert(result_string.compare(EXPECTED[i]) == 0); \ - } \ } \ } @@ -125,16 +115,8 @@ int main() { 0.0001, ma6); TEST2(sycl::remquo, float, int, 3, EXPECTED(float, 1.4f, 4.2f, 5.3f), EXPECTED(int, 0, 0, 0), 0.0001, ma6, ma3); - auto backend = deviceQueue.get_device().get_backend(); - // TODO: enable for all backends when OpenCL CPU and OpenCL/Level Zero GPU - // return correct results for nan function - if (backend == sycl::backend::ext_oneapi_cuda || - backend == sycl::backend::ext_oneapi_hip) { - TEST3(sycl::nan, float, 3, EXPECTED(std::string, "00001", "00002", "00003"), - ma7); - TEST3(sycl::nan, double, 3, - EXPECTED(std::string, "00001", "00002", "00003"), ma8); - } + TEST3(sycl::nan, float, 3, ma7); + TEST3(sycl::nan, double, 3, ma8); TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1, ma1);