diff --git a/sycl/test-e2e/Basic/built-ins/helpers.hpp b/sycl/test-e2e/Basic/built-ins/helpers.hpp index 788055369650e..2cc4fac936bc8 100644 --- a/sycl/test-e2e/Basic/built-ins/helpers.hpp +++ b/sycl/test-e2e/Basic/built-ins/helpers.hpp @@ -1,12 +1,12 @@ #include -template bool equal(T x, T y, double delta) { +template bool equal(T x, T y) { // Maybe should be C++20's std::equality_comparable. if constexpr (std::is_scalar_v) { - return std::abs(x - y) <= delta; + return x == y; } else { for (size_t i = 0; i < x.size(); ++i) - if (std::abs(x[i] - y[i]) > delta) + if (x[i] != y[i]) return false; return true; @@ -15,10 +15,10 @@ template bool equal(T x, T y, double delta) { template -void test(bool CheckDevice, double delta, FuncTy F, ExpectedTy Expected, ArgTys... Args) { +void test(bool CheckDevice, FuncTy F, ExpectedTy Expected, ArgTys... Args) { auto R = F(Args...); static_assert(std::is_same_v); - assert(equal(R, Expected, delta)); + assert(equal(R, Expected)); if (!CheckDevice) return; @@ -29,7 +29,7 @@ void test(bool CheckDevice, double delta, FuncTy F, ExpectedTy Expected, ArgTys. cgh.single_task([=]() { auto R = F(Args...); static_assert(std::is_same_v); - Success[0] = equal(R, Expected, delta); + Success[0] = equal(R, Expected); }); }); assert(sycl::host_accessor{SuccessBuf}[0]); @@ -37,16 +37,7 @@ void test(bool CheckDevice, double delta, FuncTy F, ExpectedTy Expected, ArgTys. template void test(FuncTy F, ExpectedTy Expected, ArgTys... Args) { - test(true /*CheckDevice*/, 0.0 /*delta*/, F, Expected, Args...); -} -template -void test(bool CheckDevice, FuncTy F, ExpectedTy Expected, ArgTys... Args) { - test(CheckDevice, 0.0 /*delta*/, F, Expected, Args...); -} -template -void test(double delta, FuncTy F, ExpectedTy Expected, ArgTys... Args) { - test(true /*CheckDevice*/, delta, F, Expected, Args...); + test(true /*CheckDevice*/, F, Expected, Args...); } // MSVC's STL spoils global namespace with math functions, so use explicit diff --git a/sycl/test-e2e/Basic/built-ins/marray_common.cpp b/sycl/test-e2e/Basic/built-ins/marray_common.cpp index f9415802c6ecf..2c63690f39f1c 100644 --- a/sycl/test-e2e/Basic/built-ins/marray_common.cpp +++ b/sycl/test-e2e/Basic/built-ins/marray_common.cpp @@ -8,76 +8,115 @@ #endif #include -#include "helpers.hpp" +#include -int main() { - using namespace sycl; +#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); \ + } \ + } - queue deviceQueue; - device dev = deviceQueue.get_device(); +#define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__}) - marray ma1{1.0f, 2.0f}; - marray ma2{1.0f, 2.0f}; - marray ma3{3.0f, 2.0f}; - marray ma4{1.0, 2.0}; - marray ma5{M_PI, M_PI, M_PI}; - marray ma6{M_PI, M_PI, M_PI}; - marray ma7{M_PI, M_PI, M_PI}; - marray ma8{0.3f, 0.6f}; - marray ma9{5.0, 8.0}; - marray ma10{180, 180, 180}; - marray ma11{180, 180, 180}; - marray ma12{180, 180, 180}; - marray ma13{181, 179, 181}; - marray ma14{+0.0f, -0.6f}; - marray ma15{-0.0, 0.6f}; +int main() { + sycl::queue deviceQueue; + sycl::device dev = deviceQueue.get_device(); - bool has_fp16 = queue{}.get_device().has(sycl::aspect::fp16); - bool has_fp64 = queue{}.get_device().has(sycl::aspect::fp64); + 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}; - // clamp - test(F(clamp), marray{1.0f, 2.0f}, ma1, ma2, ma3); - test(F(clamp), marray{1.0f, 2.0f}, ma1, 1.0f, 3.0f); - test(has_fp64, F(clamp), marray{1.0, 2.0}, ma4, 1.0, 3.0); - // degrees - test(F(degrees), marray{180, 180, 180}, ma5); - test(has_fp64, F(degrees), marray{180, 180, 180}, ma6); - test(has_fp16, 0.2, F(degrees), marray{180, 180, 180}, ma7); - // max - test(F(max), marray{3.0f, 2.0f}, ma1, ma3); - test(F(max), marray{1.5f, 2.0f}, ma1, 1.5f); - test(has_fp64, F(max), marray{1.5, 2.0}, ma4, 1.5); - // min - test(F(min), marray{1.0f, 2.0f}, ma1, ma3); - test(F(min), marray{1.0f, 1.5f}, ma1, 1.5f); - test(has_fp64, F(min), marray{1.0, 1.5}, ma4, 1.5); - // mix - test(F(mix), marray{1.6f, 2.0f}, ma1, ma3, ma8); - test(F(mix), marray{1.4f, 2.0f}, ma1, ma3, 0.2); - test(has_fp64, F(mix), marray{3.0, 5.0}, ma4, ma9, 0.5); - // radians - test(F(radians), marray{M_PI, M_PI, M_PI}, ma10); - test(has_fp64, F(radians), marray{M_PI, M_PI, M_PI}, ma11); - test(has_fp16, 0.002, F(radians), marray{M_PI, M_PI, M_PI}, ma12); - // step - test(F(step), marray{1.0f, 1.0f}, ma1, ma3); - test(has_fp64, F(step), marray{1.0, 1.0}, ma4, ma9); - test(has_fp16, F(step), marray{1.0, 0.0, 1.0}, ma12, ma13); - test(F(step), marray{1.0f, 0.0f}, 2.5f, ma3); - test(has_fp64, F(step), marray{0.0f, 1.0f}, 6.0f, ma9); - // smoothstep - test(F(smoothstep), marray{1.0f, 1.0f}, ma8, ma1, ma2); - test(has_fp64, 0.00000001, F(smoothstep), marray{1.0, 1.0f}, ma4, - ma9, ma9); - test(has_fp16, F(smoothstep), marray{1.0, 1.0, 1.0}, ma7, ma12, - ma13); - test(0.0000001, F(smoothstep), marray{0.0553936f, 0.0f}, 2.5f, 6.0f, - ma3); - test(has_fp64, F(smoothstep), marray{0.0f, 1.0f}, 6.0f, 8.0f, ma9); + // sycl::clamp + TEST(sycl::clamp, float, 2, EXPECTED(float, 1.0f, 2.0f), 0, ma1, ma2, ma3); + 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, 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)) + 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, 1.0f), 0, ma8, ma1, + ma2); + if (dev.has(sycl::aspect::fp64)) + 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, 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)) + TEST(sycl::smoothstep, double, 2, EXPECTED(double, 0.0f, 1.0f), 0, 6.0f, + 8.0f, ma9); // sign - test(F(sign), marray{+0.0f, -1.0f}, ma14); - test(has_fp64, F(sign), marray{-0.0, 1.0}, ma15); - test(has_fp16, F(sign), marray{1.0, 1.0, 1.0}, ma12); + 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; }