From c3a5cd5a9eea20a12a54860105edbeb7e7e83963 Mon Sep 17 00:00:00 2001 From: haonanya Date: Fri, 6 Mar 2020 11:13:20 +0800 Subject: [PATCH 01/19] [SYCL] add test for all math function. Signed-off-by: haonanya --- sycl/test/devicelib/cmath_fp64_test.cpp | 914 ++++++++++++++++++++++++ sycl/test/devicelib/cmath_test.cpp | 842 +++++++++++++++++++++- sycl/test/devicelib/cmath_test_fp64.cpp | 86 --- sycl/test/devicelib/math_fp64_test.cpp | 798 ++++++++++++++++++++- sycl/test/devicelib/math_test.cpp | 799 ++++++++++++++++++++- 5 files changed, 3304 insertions(+), 135 deletions(-) create mode 100644 sycl/test/devicelib/cmath_fp64_test.cpp delete mode 100644 sycl/test/devicelib/cmath_test_fp64.cpp diff --git a/sycl/test/devicelib/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp new file mode 100644 index 0000000000000..279a75c29658f --- /dev/null +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -0,0 +1,914 @@ +// UNSUPPORTED: windows +// RUN: %clangxx -fsycl -c %s -o %t.o +// RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath-fp64.o -o %t.out +#include +#include +#include + +namespace s = cl::sycl; +constexpr s::access::mode sycl_read = s::access::mode::read; +constexpr s::access::mode sycl_write = s::access::mode::write; + +template +class DeviceCos; + +template +void device_cos_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::cos(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceSin; + +template +void device_sin_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::sin(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLog; + +template +void device_log_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::log(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAcos; + +template +void device_acos_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::acos(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAsin; + +template +void device_asin_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::asin(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAtan; + +template +void device_atan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::atan(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAtan2; + +template +void device_atan2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + T b = 1; + res_access[0] = std::atan2(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceCosh; + +template +void device_cosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::cosh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceExp; + +template +void device_exp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::exp(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFmod; + +template +void device_fmod_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1.5; + T b = 1; + res_access[0] = std::fmod(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFrexp; + +template +void device_frexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + int exp = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&exp, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto exp_access = buffer2.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::frexp(a, &exp_access[0]); + }); + }); + } + + assert(result == ref && exp == 0); +} + +template +class DeviceLdexp; + +template +void device_ldexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::ldexp(a, 1); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLog10; + +template +void device_log10_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::log10(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceModf; + +template +void device_modf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T iptr = -1; + T ref1 = 0; + T ref2 = 1; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&iptr, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto iptr_access = buffer2.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::modf(a, &iptr_access[0]); + }); + }); + } + + assert(result == ref1 && iptr == ref2); +} + +template +class DevicePow; + +template +void device_pow_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + T b = 1; + res_access[0] = std::pow(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceSinh; + +template +void device_sinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::sinh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceSqrt; + +template +void device_sqrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 4; + res_access[0] = std::sqrt(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceTan; + +template +void device_tan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::tan(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceTanh; + +template +void device_tanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::tanh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAcosh; + +template +void device_acosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::acosh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAsinh; + +template +void device_asinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::asinh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAtanh; + +template +void device_atanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::atanh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceCbrt; + +template +void device_cbrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::cbrt(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceErf; + +template +void device_erf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::erf(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceErfc; + +template +void device_erfc_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::erfc(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceExp2; + +template +void device_exp2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::exp2(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceExpm1; + +template +void device_expm1_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::expm1(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFdim; + +template +void device_fdim_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + T b = 0; + res_access[0] = std::fdim(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFma; + +template +void device_fma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + T b = 1; + T c = 1; + res_access[0] = std::fma(a, b, c); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceHypot; + +template +void device_hypot_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 3; + T b = 4; + res_access[0] = std::hypot(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceIlogb; + +template +void device_ilogb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::ilogb(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLgamma; + +template +void device_lgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = NAN; + res_access[0] = std::lgamma(a); + }); + }); + } + + assert(std::isnan(result)); +} + +template +class DeviceLog1p; + +template +void device_log1p_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::log1p(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLog2; + +template +void device_log2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::log2(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLogb; + +template +void device_logb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::logb(1); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceRemainder; + +template +void device_remainder_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0.5; + T b = 1; + res_access[0] = std::remainder(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceRemquo; + +template +void device_remquo_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + int quo = -1; + T ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&quo, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto quo_access = buffer2.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0.5; + T b = 1; + res_access[0] = std::remquo(a, b, &quo_access[0]); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceTgamma; + +template +void device_tgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = NAN; + res_access[0] = std::tgamma(a); + }); + }); + } + + assert(std::isnan(result)); +} + +template +void device_cmath_test(s::queue &deviceQueue) { + device_cos_test(deviceQueue); + device_sin_test(deviceQueue); + device_log_test(deviceQueue); + device_acos_test(deviceQueue); + device_asin_test(deviceQueue); + device_atan_test(deviceQueue); + device_atan2_test(deviceQueue); + device_cosh_test(deviceQueue); + device_exp_test(deviceQueue); + device_fmod_test(deviceQueue); + device_frexp_test(deviceQueue); + device_ldexp_test(deviceQueue); + device_log10_test(deviceQueue); + device_modf_test(deviceQueue); + device_pow_test(deviceQueue); + device_sinh_test(deviceQueue); + device_sqrt_test(deviceQueue); + device_tan_test(deviceQueue); + device_tanh_test(deviceQueue); + device_acosh_test(deviceQueue); + device_asinh_test(deviceQueue); + device_atanh_test(deviceQueue); + device_cbrt_test(deviceQueue); + device_erf_test(deviceQueue); + device_erfc_test(deviceQueue); + device_exp2_test(deviceQueue); + device_expm1_test(deviceQueue); + device_fdim_test(deviceQueue); + device_fma_test(deviceQueue); + device_hypot_test(deviceQueue); + device_ilogb_test(deviceQueue); + device_lgamma_test(deviceQueue); + device_log1p_test(deviceQueue); + device_log2_test(deviceQueue); + device_logb_test(deviceQueue); + device_remainder_test(deviceQueue); + device_remquo_test(deviceQueue); + device_tgamma_test(deviceQueue); +} + +int main() { + s::queue deviceQueue; + if (deviceQueue.get_device().has_extension("cl_khr_fp64")) { + device_cmath_test(deviceQueue); + std::cout << "Pass" << std::endl; + } + return 0; +} diff --git a/sycl/test/devicelib/cmath_test.cpp b/sycl/test/devicelib/cmath_test.cpp index 217ad4121f6c8..62ae26b2217c1 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -2,8 +2,8 @@ // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath.o -o %t.out #include -#include #include +#include namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; @@ -16,17 +16,19 @@ template void device_cos_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; T result = -1; + T ref = 1; { s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); cgh.single_task >([=]() { - res_access[0] = std::cos(0); + T a = 0; + res_access[0] = std::cos(a); }); }); } - assert(result == 1); + assert(result == ref); } template @@ -36,17 +38,19 @@ template void device_sin_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; T result = -1; + T ref = 0; { s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); cgh.single_task >([=]() { - res_access[0] = std::sin(0); + T a = 0; + res_access[0] = std::sin(a); }); }); } - assert(result == 0); + assert(result == ref); } template @@ -56,17 +60,806 @@ template void device_log_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; T result = -1; + T ref = 0; { s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); cgh.single_task >([=]() { - res_access[0] = std::log(1); + T a = 1; + res_access[0] = std::log(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAcos; + +template +void device_acos_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::acos(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAsin; + +template +void device_asin_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::asin(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAtan; + +template +void device_atan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::atan(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAtan2; + +template +void device_atan2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + T b = 1; + res_access[0] = std::atan2(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceCosh; + +template +void device_cosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::cosh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceExp; + +template +void device_exp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::exp(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFmod; + +template +void device_fmod_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1.5; + T b = 1; + res_access[0] = std::fmod(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFrexp; + +template +void device_frexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + int exp = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&exp, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto exp_access = buffer2.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::frexp(a, &exp_access[0]); + }); + }); + } + + assert(result == ref && exp == 0); +} + +template +class DeviceLdexp; + +template +void device_ldexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::ldexp(a, 1); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLog10; + +template +void device_log10_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::log10(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceModf; + +template +void device_modf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T iptr = -1; + T ref1 = 0; + T ref2 = 1; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&iptr, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto iptr_access = buffer2.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::modf(a, &iptr_access[0]); + }); + }); + } + + assert(result == ref1 && iptr == ref2); +} + +template +class DevicePow; + +template +void device_pow_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + T b = 1; + res_access[0] = std::pow(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceSinh; + +template +void device_sinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::sinh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceSqrt; + +template +void device_sqrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 4; + res_access[0] = std::sqrt(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceTan; + +template +void device_tan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::tan(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceTanh; + +template +void device_tanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::tanh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAcosh; + +template +void device_acosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::acosh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAsinh; + +template +void device_asinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::asinh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceAtanh; + +template +void device_atanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::atanh(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceCbrt; + +template +void device_cbrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::cbrt(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceErf; + +template +void device_erf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::erf(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceErfc; + +template +void device_erfc_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::erfc(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceExp2; + +template +void device_exp2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::exp2(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceExpm1; + +template +void device_expm1_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::expm1(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFdim; + +template +void device_fdim_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + T b = 0; + res_access[0] = std::fdim(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceFma; + +template +void device_fma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + T b = 1; + T c = 1; + res_access[0] = std::fma(a, b, c); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceHypot; + +template +void device_hypot_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 3; + T b = 4; + res_access[0] = std::hypot(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceIlogb; + +template +void device_ilogb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::ilogb(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLgamma; + +template +void device_lgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = NAN; + res_access[0] = std::lgamma(a); + }); + }); + } + + assert(std::isnan(result)); +} + +template +class DeviceLog1p; + +template +void device_log1p_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0; + res_access[0] = std::log1p(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLog2; + +template +void device_log2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::log2(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceLogb; + +template +void device_logb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 1; + res_access[0] = std::logb(a); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceRemainder; + +template +void device_remainder_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + T ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0.5; + T b = 1; + res_access[0] = std::remainder(a, b); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceRemquo; + +template +void device_remquo_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + int quo = -1; + T ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&quo, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto quo_access = buffer2.template get_access(cgh); + cgh.single_task >([=]() { + T a = 0.5; + T b = 1; + res_access[0] = std::remquo(a, b, &quo_access[0]); + }); + }); + } + + assert(result == ref); +} + +template +class DeviceTgamma; + +template +void device_tgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + T result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task >([=]() { + T a = NAN; + res_access[0] = std::tgamma(a); }); }); } - assert(result == 0); + assert(std::isnan(result)); } template @@ -74,6 +867,41 @@ void device_cmath_test(s::queue &deviceQueue) { device_cos_test(deviceQueue); device_sin_test(deviceQueue); device_log_test(deviceQueue); + device_acos_test(deviceQueue); + device_asin_test(deviceQueue); + device_atan_test(deviceQueue); + device_atan2_test(deviceQueue); + device_cosh_test(deviceQueue); + device_exp_test(deviceQueue); + device_fmod_test(deviceQueue); + device_frexp_test(deviceQueue); + device_ldexp_test(deviceQueue); + device_log10_test(deviceQueue); + device_modf_test(deviceQueue); + device_pow_test(deviceQueue); + device_sinh_test(deviceQueue); + device_sqrt_test(deviceQueue); + device_tan_test(deviceQueue); + device_tanh_test(deviceQueue); + device_acosh_test(deviceQueue); + device_asinh_test(deviceQueue); + device_atanh_test(deviceQueue); + device_cbrt_test(deviceQueue); + device_erf_test(deviceQueue); + device_erfc_test(deviceQueue); + device_exp2_test(deviceQueue); + device_expm1_test(deviceQueue); + device_fdim_test(deviceQueue); + device_fma_test(deviceQueue); + device_hypot_test(deviceQueue); + device_ilogb_test(deviceQueue); + device_lgamma_test(deviceQueue); + device_log1p_test(deviceQueue); + device_log2_test(deviceQueue); + device_logb_test(deviceQueue); + device_remainder_test(deviceQueue); + device_remquo_test(deviceQueue); + device_tgamma_test(deviceQueue); } int main() { diff --git a/sycl/test/devicelib/cmath_test_fp64.cpp b/sycl/test/devicelib/cmath_test_fp64.cpp deleted file mode 100644 index 1c8b7afa5d4a5..0000000000000 --- a/sycl/test/devicelib/cmath_test_fp64.cpp +++ /dev/null @@ -1,86 +0,0 @@ -// UNSUPPORTED: windows -// RUN: %clangxx -fsycl -c %s -o %t.o -// RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath-fp64.o -o %t.out -#include -#include -#include - -namespace s = cl::sycl; -constexpr s::access::mode sycl_read = s::access::mode::read; -constexpr s::access::mode sycl_write = s::access::mode::write; - -template -class DeviceCos; - -template -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - res_access[0] = std::cos(0); - }); - }); - } - - assert(result == 1); -} - -template -class DeviceSin; - -template -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - res_access[0] = std::sin(0); - }); - }); - } - - assert(result == 0); -} - -template -class DeviceLog; - -template -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - res_access[0] = std::log(1); - }); - }); - } - - assert(result == 0); -} - -template -void device_cmath_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); -} - -int main() { - s::queue deviceQueue; - if (deviceQueue.get_device().has_extension("cl_khr_fp64")) { - device_cmath_test(deviceQueue); - std::cout << "Pass" << std::endl; - } - return 0; -} diff --git a/sycl/test/devicelib/math_fp64_test.cpp b/sycl/test/devicelib/math_fp64_test.cpp index 30c21dbc3b77e..a56da6dea2301 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -9,64 +9,820 @@ namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -class DeviceSin; +class DeviceCos; -void device_sin_test(s::queue &deviceQueue) { +void device_cos_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; - double result_d = -1; + double result = -1; + double ref = 1; { - s::buffer buffer1(&result_d, numOfItems); + s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = sin(0); + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = cos(a); }); }); } - assert(result_d == 0); + assert(result == ref); } -class DeviceCos; +class DeviceSin; -void device_cos_test(s::queue &deviceQueue) { +void device_sin_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; - double result_d = -1; + double result = -1; + double ref = 0; { - s::buffer buffer1(&result_d, numOfItems); + s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = cos(0); + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = sin(a); }); }); } - assert(result_d == 1); + assert(result == ref); } class DeviceLog; void device_log_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; - double result_d = -1; + double result = -1; + double ref = 0; { - s::buffer buffer2(&result_d, numOfItems); + s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer2.get_access(cgh); + auto res_access = buffer1.template get_access(cgh); cgh.single_task([=]() { - res_access1[0] = log(1); + double a = 1; + res_access[0] = log(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAcos; + +void device_acos_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = acos(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAsin; + +void device_asin_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = asin(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAtan; + +void device_atan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = atan(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAtan2; + +void device_atan2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + double b = 1; + res_access[0] = atan2(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceCosh; + +void device_cosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = cosh(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceExp; + +void device_exp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = exp(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceFmod; + +void device_fmod_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1.5; + double b = 1; + res_access[0] = fmod(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceFrexp; + +void device_frexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + int exp = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&exp, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto exp_access = buffer2.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = frexp(a, &exp_access[0]); + }); + }); + } + + assert(result == ref && exp == 0); +} + +class DeviceLdexp; + +void device_ldexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = ldexp(a, 1); + }); + }); + } + + assert(result == ref); +} + +class DeviceLog10; + +void device_log10_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = log10(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceModf; + +void device_modf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double iptr = -1; + double ref1 = 0; + double ref2 = 1; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&iptr, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto iptr_access = buffer2.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = modf(a, &iptr_access[0]); + }); + }); + } + + assert(result == ref1 && iptr == ref2); +} + +class DevicePow; + +void device_pow_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + double b = 1; + res_access[0] = pow(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceSinh; + +void device_sinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = sinh(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceSqrt; + +void device_sqrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 4; + res_access[0] = sqrt(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceTan; + +void device_tan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = tan(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceTanh; + +void device_tanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = tanh(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAcosh; + +void device_acosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = acosh(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAsinh; + +void device_asinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = asinh(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAtanh; + +void device_atanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = atanh(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceCbrt; + +void device_cbrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = cbrt(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceErf; + +void device_erf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = erf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceErfc; + +void device_erfc_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = erfc(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceExp2; + +void device_exp2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = exp2(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceExpm1; + +void device_expm1_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = expm1(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceFdim; + +void device_fdim_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + double b = 0; + res_access[0] = fdim(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceFma; + +void device_fma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + double b = 1; + double c = 1; + res_access[0] = fma(a, b, c); + }); + }); + } + + assert(result == ref); +} + +class DeviceHypot; + +void device_hypot_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 3; + double b = 4; + res_access[0] = hypot(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceIlogb; + +void device_ilogb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = ilogb(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceLgamma; + +void device_lgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = NAN; + res_access[0] = lgamma(a); + }); + }); + } + + assert(std::isnan(result)); +} + +class DeviceLog1p; + +void device_log1p_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0; + res_access[0] = log1p(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceLog2; + +void device_log2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = log2(a); }); }); } - assert(result_d == 0); + assert(result == ref); } +class DeviceLogb; + +void device_logb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 1; + res_access[0] = logb(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceRemainder; + +void device_remainder_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + double ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = 0.5; + double b = 1; + res_access[0] = remainder(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceRemquo; + +void device_remquo_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + int quo = -1; + double ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&quo, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto quo_access = buffer2.template get_access(cgh); + cgh.single_task([=]() { + double a = 0.5; + double b = 1; + res_access[0] = remquo(a, b, &quo_access[0]); + }); + }); + } + + assert(result == ref); +} + +class DeviceTgamma; +void device_tgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + double result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + double a = NAN; + res_access[0] = tgamma(a); + }); + }); + } + + assert(std::isnan(result)); +} void device_math_test(s::queue &deviceQueue) { device_cos_test(deviceQueue); device_sin_test(deviceQueue); device_log_test(deviceQueue); + device_acos_test(deviceQueue); + device_asin_test(deviceQueue); + device_atan_test(deviceQueue); + device_atan2_test(deviceQueue); + device_cosh_test(deviceQueue); + device_exp_test(deviceQueue); + device_fmod_test(deviceQueue); + device_frexp_test(deviceQueue); + device_ldexp_test(deviceQueue); + device_log10_test(deviceQueue); + device_modf_test(deviceQueue); + device_pow_test(deviceQueue); + device_sinh_test(deviceQueue); + device_sqrt_test(deviceQueue); + device_tan_test(deviceQueue); + device_tanh_test(deviceQueue); + device_acosh_test(deviceQueue); + device_asinh_test(deviceQueue); + device_atanh_test(deviceQueue); + device_cbrt_test(deviceQueue); + device_erf_test(deviceQueue); + device_erfc_test(deviceQueue); + device_exp2_test(deviceQueue); + device_expm1_test(deviceQueue); + device_fdim_test(deviceQueue); + device_fma_test(deviceQueue); + device_hypot_test(deviceQueue); + device_ilogb_test(deviceQueue); + device_lgamma_test(deviceQueue); + device_log1p_test(deviceQueue); + device_log2_test(deviceQueue); + device_logb_test(deviceQueue); + device_remainder_test(deviceQueue); + device_remquo_test(deviceQueue); + device_tgamma_test(deviceQueue); } int main() { diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index 1b5aa4332fe8e..f60f376b52443 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -9,64 +9,821 @@ namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -class DeviceSin; +class DeviceCos; -void device_sin_test(s::queue &deviceQueue) { +void device_cos_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; - float result_f = -1; + float result = -1; + float ref = 1; { - s::buffer buffer1(&result_f, numOfItems); + s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = sinf(0); + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = cosf(a); }); }); } - assert(result_f == 0); + assert(result == ref); } -class DeviceCos; +class DeviceSin; -void device_cos_test(s::queue &deviceQueue) { +void device_sin_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; - float result_f = -1; + float result = -1; + float ref = 0; { - s::buffer buffer1(&result_f, numOfItems); + s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = cosf(0); + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = sinf(a); }); }); } - assert(result_f == 1); + assert(result == ref); } class DeviceLog; void device_log_test(s::queue &deviceQueue) { s::range<1> numOfItems{1}; - float result_f = -1; + float result = -1; + float ref = 0; { - s::buffer buffer1(&result_f, numOfItems); + s::buffer buffer1(&result, numOfItems); deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); + auto res_access = buffer1.template get_access(cgh); cgh.single_task([=]() { - res_access1[0] = logf(1); + float a = 1; + res_access[0] = logf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAcos; + +void device_acos_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = acosf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAsin; + +void device_asin_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = asinf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAtan; + +void device_atan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = atanf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAtan2; + +void device_atan2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + float b = 1; + res_access[0] = atan2f(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceCosh; + +void device_cosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = coshf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceExp; + +void device_exp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = expf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceFmod; + +void device_fmod_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1.5; + float b = 1; + res_access[0] = fmodf(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceFrexp; + +void device_frexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + int exp = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&exp, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto exp_access = buffer2.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = frexpf(a, &exp_access[0]); + }); + }); + } + + assert(result == ref && exp == 0); +} + +class DeviceLdexp; + +void device_ldexp_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = ldexpf(a, 1); + }); + }); + } + + assert(result == ref); +} + +class DeviceLog10; + +void device_log10_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = log10f(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceModf; + +void device_modf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float iptr = -1; + float ref1 = 0; + float ref2 = 1; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&iptr, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto iptr_access = buffer2.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = modff(a, &iptr_access[0]); + }); + }); + } + + assert(result == ref1 && iptr == ref2); +} + +class DevicePow; + +void device_pow_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + float b = 1; + res_access[0] = powf(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceSinh; + +void device_sinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = sinhf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceSqrt; + +void device_sqrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 4; + res_access[0] = sqrtf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceTan; + +void device_tan_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = tanf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceTanh; + +void device_tanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = tanhf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAcosh; + +void device_acosh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = acoshf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAsinh; + +void device_asinh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = asinhf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceAtanh; + +void device_atanh_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = atanhf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceCbrt; + +void device_cbrt_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = cbrtf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceErf; + +void device_erf_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = erff(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceErfc; + +void device_erfc_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = erfcf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceExp2; + +void device_exp2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = exp2f(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceExpm1; + +void device_expm1_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = expm1f(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceFdim; + +void device_fdim_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = fdimf(1, a); + }); + }); + } + + assert(result == ref); +} + +class DeviceFma; + +void device_fma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 2; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + float b = 1; + float c = 1; + res_access[0] = fmaf(a, b, c); + }); + }); + } + + assert(result == ref); +} + +class DeviceHypot; + +void device_hypot_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 3; + float b = 4; + res_access[0] = hypotf(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceIlogb; + +void device_ilogb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = ilogbf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceLgamma; + +void device_lgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = NAN; + res_access[0] = lgammaf(a); + }); + }); + } + + assert(std::isnan(result)); +} + +class DeviceLog1p; + +void device_log1p_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0; + res_access[0] = log1pf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceLog2; + +void device_log2_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = log2f(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceLogb; + +void device_logb_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 1; + res_access[0] = logbf(a); + }); + }); + } + + assert(result == ref); +} + +class DeviceRemainder; + +void device_remainder_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + float ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = 0.5; + float b = 1; + res_access[0] = remainderf(a, b); + }); + }); + } + + assert(result == ref); +} + +class DeviceRemquo; + +void device_remquo_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + int quo = -1; + float ref = 0.5; + { + s::buffer buffer1(&result, numOfItems); + s::buffer buffer2(&quo, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + auto quo_access = buffer2.template get_access(cgh); + cgh.single_task([=]() { + float a = 0.5; + float b = 1; + res_access[0] = remquof(a, b, &quo_access[0]); + }); + }); + } + + assert(result == ref); +} + +class DeviceTgamma; + +void device_tgamma_test(s::queue &deviceQueue) { + s::range<1> numOfItems{1}; + float result = -1; + { + s::buffer buffer1(&result, numOfItems); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto res_access = buffer1.template get_access(cgh); + cgh.single_task([=]() { + float a = NAN; + res_access[0] = tgammaf(a); }); }); } - assert(result_f == 0); + assert(std::isnan(result)); } void device_math_test(s::queue &deviceQueue) { device_cos_test(deviceQueue); device_sin_test(deviceQueue); device_log_test(deviceQueue); + device_acos_test(deviceQueue); + device_asin_test(deviceQueue); + device_atan_test(deviceQueue); + device_atan2_test(deviceQueue); + device_cosh_test(deviceQueue); + device_exp_test(deviceQueue); + device_fmod_test(deviceQueue); + device_frexp_test(deviceQueue); + device_ldexp_test(deviceQueue); + device_log10_test(deviceQueue); + device_modf_test(deviceQueue); + device_pow_test(deviceQueue); + device_sinh_test(deviceQueue); + device_sqrt_test(deviceQueue); + device_tan_test(deviceQueue); + device_tanh_test(deviceQueue); + device_acosh_test(deviceQueue); + device_asinh_test(deviceQueue); + device_atanh_test(deviceQueue); + device_cbrt_test(deviceQueue); + device_erf_test(deviceQueue); + device_erfc_test(deviceQueue); + device_exp2_test(deviceQueue); + device_expm1_test(deviceQueue); + device_fdim_test(deviceQueue); + device_fma_test(deviceQueue); + device_hypot_test(deviceQueue); + device_ilogb_test(deviceQueue); + device_lgamma_test(deviceQueue); + device_log1p_test(deviceQueue); + device_log2_test(deviceQueue); + device_logb_test(deviceQueue); + device_remainder_test(deviceQueue); + device_remquo_test(deviceQueue); + device_tgamma_test(deviceQueue); } int main() { From 1816430f3100947d6293a7fddaaba2185b139f5f Mon Sep 17 00:00:00 2001 From: haonanya Date: Fri, 6 Mar 2020 11:15:07 +0800 Subject: [PATCH 02/19] [SYCL] in llvm/sycl/source/detail/devicelib/CMakeLists.txt, libsycl-cmath.o depends on device_math.h, modify here Signed-off-by: haonanya --- sycl/source/detail/devicelib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/devicelib/CMakeLists.txt b/sycl/source/detail/devicelib/CMakeLists.txt index 99aab2e4c7f1e..83c523abe1518 100644 --- a/sycl/source/detail/devicelib/CMakeLists.txt +++ b/sycl/source/detail/devicelib/CMakeLists.txt @@ -64,7 +64,7 @@ add_custom_command(OUTPUT ${devicelib-obj-cmath} ${CMAKE_CURRENT_SOURCE_DIR}/cmath_wrapper.cpp -o ${devicelib-obj-cmath} MAIN_DEPENDENCY cmath_wrapper.cpp - DEPENDS device_complex.h clang + DEPENDS device_math.h clang VERBATIM) set(devicelib-obj-cmath-fp64 ${binary_dir}/libsycl-cmath-fp64.o) From 5540691d0f48b1a80ffcb8932312425355f340b8 Mon Sep 17 00:00:00 2001 From: haonanya Date: Mon, 9 Mar 2020 20:33:01 +0800 Subject: [PATCH 03/19] [SYCL] Put all math function calls into a single kernel. Signed-off-by: haonanya --- sycl/test/devicelib/cmath_fp64_test.cpp | 1116 +++++------------------ sycl/test/devicelib/cmath_test.cpp | 1116 +++++------------------ sycl/test/devicelib/math_fp64_test.cpp | 1038 +++++---------------- sycl/test/devicelib/math_test.cpp | 1039 +++++---------------- 4 files changed, 920 insertions(+), 3389 deletions(-) diff --git a/sycl/test/devicelib/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp index 279a75c29658f..bbe6a39cf819b 100644 --- a/sycl/test/devicelib/cmath_fp64_test.cpp +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -4,904 +4,248 @@ #include #include #include +#include "math_utils.hpp" namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -template -class DeviceCos; - -template -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::cos(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceSin; - -template -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::sin(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLog; - -template -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::log(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAcos; - -template -void device_acos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::acos(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAsin; - -template -void device_asin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::asin(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAtan; - -template -void device_atan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::atan(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAtan2; - -template -void device_atan2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - T b = 1; - res_access[0] = std::atan2(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceCosh; - -template -void device_cosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::cosh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceExp; - -template -void device_exp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::exp(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFmod; - -template -void device_fmod_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1.5; - T b = 1; - res_access[0] = std::fmod(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFrexp; +double a = 0; +double b = 1; +double c = 0.5; +double d = 2; +double e = 5; template -void device_frexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - int exp = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&exp, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto exp_access = buffer2.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::frexp(a, &exp_access[0]); - }); - }); - } - - assert(result == ref && exp == 0); -} - -template -class DeviceLdexp; - -template -void device_ldexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::ldexp(a, 1); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLog10; - -template -void device_log10_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::log10(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceModf; - -template -void device_modf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; +void device_cmath_test(s::queue &deviceQueue) { + s::range<1> numOfItems{38}; + T result[38] = {-1}; + T ref[38] = { + b, + a, + a, + a, + a, + a, + a, + b, + b, + c, + a, + d, + a, + a, + b, + a, + d, + a, + a, + a, + a, + a, + b, + a, + b, + d, + a, + b, + d, + e, + a, + a, + a, + a, + c, + c, + a, + a, + }; + int expv = -1; T iptr = -1; - T ref1 = 0; - T ref2 = 1; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&iptr, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto iptr_access = buffer2.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::modf(a, &iptr_access[0]); - }); - }); - } - - assert(result == ref1 && iptr == ref2); -} - -template -class DevicePow; - -template -void device_pow_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - T b = 1; - res_access[0] = std::pow(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceSinh; - -template -void device_sinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::sinh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceSqrt; - -template -void device_sqrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 4; - res_access[0] = std::sqrt(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceTan; - -template -void device_tan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::tan(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceTanh; - -template -void device_tanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::tanh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAcosh; - -template -void device_acosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::acosh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAsinh; - -template -void device_asinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::asinh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAtanh; - -template -void device_atanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::atanh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceCbrt; - -template -void device_cbrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::cbrt(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceErf; - -template -void device_erf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::erf(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceErfc; - -template -void device_erfc_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::erfc(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceExp2; - -template -void device_exp2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::exp2(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceExpm1; - -template -void device_expm1_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::expm1(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFdim; - -template -void device_fdim_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - T b = 0; - res_access[0] = std::fdim(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFma; - -template -void device_fma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - T b = 1; - T c = 1; - res_access[0] = std::fma(a, b, c); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceHypot; - -template -void device_hypot_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 3; - T b = 4; - res_access[0] = std::hypot(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceIlogb; - -template -void device_ilogb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::ilogb(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLgamma; - -template -void device_lgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = NAN; - res_access[0] = std::lgamma(a); - }); - }); - } - - assert(std::isnan(result)); -} - -template -class DeviceLog1p; - -template -void device_log1p_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::log1p(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLog2; - -template -void device_log2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::log2(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLogb; - -template -void device_logb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::logb(1); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceRemainder; - -template -void device_remainder_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0.5; - T b = 1; - res_access[0] = std::remainder(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceRemquo; - -template -void device_remquo_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; int quo = -1; - T ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&quo, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto quo_access = buffer2.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0.5; - T b = 1; - res_access[0] = std::remquo(a, b, &quo_access[0]); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceTgamma; - -template -void device_tgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; { - s::buffer buffer1(&result, numOfItems); + s::buffer buffer1(result, numOfItems); + s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer3(&iptr, s::range<1>{1}); + s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = NAN; - res_access[0] = std::tgamma(a); - }); - }); - } - - assert(std::isnan(result)); -} - -template -void device_cmath_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); - device_acos_test(deviceQueue); - device_asin_test(deviceQueue); - device_atan_test(deviceQueue); - device_atan2_test(deviceQueue); - device_cosh_test(deviceQueue); - device_exp_test(deviceQueue); - device_fmod_test(deviceQueue); - device_frexp_test(deviceQueue); - device_ldexp_test(deviceQueue); - device_log10_test(deviceQueue); - device_modf_test(deviceQueue); - device_pow_test(deviceQueue); - device_sinh_test(deviceQueue); - device_sqrt_test(deviceQueue); - device_tan_test(deviceQueue); - device_tanh_test(deviceQueue); - device_acosh_test(deviceQueue); - device_asinh_test(deviceQueue); - device_atanh_test(deviceQueue); - device_cbrt_test(deviceQueue); - device_erf_test(deviceQueue); - device_erfc_test(deviceQueue); - device_exp2_test(deviceQueue); - device_expm1_test(deviceQueue); - device_fdim_test(deviceQueue); - device_fma_test(deviceQueue); - device_hypot_test(deviceQueue); - device_ilogb_test(deviceQueue); - device_lgamma_test(deviceQueue); - device_log1p_test(deviceQueue); - device_log2_test(deviceQueue); - device_logb_test(deviceQueue); - device_remainder_test(deviceQueue); - device_remquo_test(deviceQueue); - device_tgamma_test(deviceQueue); + auto exp_access = buffer2.template get_access(cgh); + auto iptr_access = buffer3.template get_access(cgh); + auto quo_access = buffer4.template get_access(cgh); + cgh.single_task([=]() { + int i = 0; + { + T a = 0; + res_access[i++] = std::cos(a); + } + { + T a = 0; + res_access[i++] = std::sin(a); + } + { + T a = 1; + res_access[i++] = std::log(a); + } + { + T a = 1; + res_access[i++] = std::acos(a); + } + { + T a = 0; + res_access[i++] = std::asin(a); + } + { + T a = 0; + res_access[i++] = std::atan(a); + } + { + T a = 0; + T b = 1; + res_access[i++] = std::atan2(a, b); + } + { + T a = 0; + res_access[i++] = std::cosh(a); + } + { + T a = 0; + res_access[i++] = std::exp(a); + } + { + T a = 1.5; + T b = 1; + res_access[i++] = std::fmod(a, b); + } + { + T a = 0; + res_access[i++] = std::frexp(a, &exp_access[0]); + } + { + T a = 1; + res_access[i++] = std::ldexp(a, 1); + } + { + T a = 1; + res_access[i++] = std::log10(a); + } + { + T a = 1; + res_access[i++] = std::modf(a, &iptr_access[0]); + } + { + T a = 1; + T b = 1; + res_access[i++] = std::pow(a, b); + } + { + T a = 0; + res_access[i++] = std::sinh(a); + } + { + T a = 4; + res_access[i++] = std::sqrt(a); + } + { + T a = 0; + res_access[i++] = std::tan(a); + } + { + T a = 0; + res_access[i++] = std::tanh(a); + } + { + T a = 1; + res_access[i++] = std::acosh(a); + } + { + T a = 0; + res_access[i++] = std::asinh(a); + } + { + T a = 0; + res_access[i++] = std::atanh(a); + } + { + T a = 1; + res_access[i++] = std::cbrt(a); + } + { + T a = 0; + res_access[i++] = std::erf(a); + } + { + T a = 0; + res_access[i++] = std::erfc(a); + } + { + T a = 1; + res_access[i++] = std::exp2(a); + } + { + T a = 0; + res_access[i++] = std::expm1(a); + } + { + T a = 0; + res_access[i++] = std::fdim(1, a); + } + { + T a = 1; + T b = 1; + T c = 1; + res_access[i++] = std::fma(a, b, c); + } + { + T a = 3; + T b = 4; + res_access[i++] = std::hypot(a, b); + } + { + T a = 1; + res_access[i++] = std::ilogb(a); + } + { + T a = 0; + res_access[i++] = std::log1p(a); + } + { + T a = 1; + res_access[i++] = std::log2(a); + } + { + T a = 1; + res_access[i++] = std::logb(a); + } + { + T a = 0.5; + T b = 1; + res_access[i++] = std::remainder(a, b); + } + { + T a = 0.5; + T b = 1; + res_access[i++] = std::remquo(a, b, &quo_access[0]); + } + { + T a = NAN; + res_access[i++] = std::tgamma(a); + } + { + T a = NAN; + res_access[i++] = std::lgamma(a); + } + }); + }); + } + for (int i = 0; i < 36; ++i) { + assert(is_about_FP(result[i], ref[i])); + } + assert(std::isnan(result[36])); + assert(std::isnan(result[37])); + assert(is_about_FP(iptr, b)); + assert(expv == 0); + assert(quo == 0); } int main() { diff --git a/sycl/test/devicelib/cmath_test.cpp b/sycl/test/devicelib/cmath_test.cpp index 62ae26b2217c1..7f6866e0a8b61 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -4,904 +4,248 @@ #include #include #include +#include "math_utils.hpp" namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -template -class DeviceCos; - -template -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::cos(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceSin; - -template -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::sin(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLog; - -template -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::log(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAcos; - -template -void device_acos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::acos(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAsin; - -template -void device_asin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::asin(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAtan; - -template -void device_atan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::atan(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAtan2; - -template -void device_atan2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - T b = 1; - res_access[0] = std::atan2(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceCosh; - -template -void device_cosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::cosh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceExp; - -template -void device_exp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::exp(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFmod; - -template -void device_fmod_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1.5; - T b = 1; - res_access[0] = std::fmod(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFrexp; +float a = 0; +float b = 1; +float c = 0.5; +float d = 2; +float e = 5; template -void device_frexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - int exp = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&exp, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto exp_access = buffer2.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::frexp(a, &exp_access[0]); - }); - }); - } - - assert(result == ref && exp == 0); -} - -template -class DeviceLdexp; - -template -void device_ldexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::ldexp(a, 1); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLog10; - -template -void device_log10_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::log10(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceModf; - -template -void device_modf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; +void device_cmath_test(s::queue &deviceQueue) { + s::range<1> numOfItems{38}; + T result[38] = {-1}; + T ref[38] = { + b, + a, + a, + a, + a, + a, + a, + b, + b, + c, + a, + d, + a, + a, + b, + a, + d, + a, + a, + a, + a, + a, + b, + a, + b, + d, + a, + b, + d, + e, + a, + a, + a, + a, + c, + c, + a, + a, + }; + int expv = -1; T iptr = -1; - T ref1 = 0; - T ref2 = 1; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&iptr, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto iptr_access = buffer2.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::modf(a, &iptr_access[0]); - }); - }); - } - - assert(result == ref1 && iptr == ref2); -} - -template -class DevicePow; - -template -void device_pow_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - T b = 1; - res_access[0] = std::pow(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceSinh; - -template -void device_sinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::sinh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceSqrt; - -template -void device_sqrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 4; - res_access[0] = std::sqrt(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceTan; - -template -void device_tan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::tan(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceTanh; - -template -void device_tanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::tanh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAcosh; - -template -void device_acosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::acosh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAsinh; - -template -void device_asinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::asinh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceAtanh; - -template -void device_atanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::atanh(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceCbrt; - -template -void device_cbrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::cbrt(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceErf; - -template -void device_erf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::erf(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceErfc; - -template -void device_erfc_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::erfc(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceExp2; - -template -void device_exp2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::exp2(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceExpm1; - -template -void device_expm1_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::expm1(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFdim; - -template -void device_fdim_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - T b = 0; - res_access[0] = std::fdim(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceFma; - -template -void device_fma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - T b = 1; - T c = 1; - res_access[0] = std::fma(a, b, c); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceHypot; - -template -void device_hypot_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 3; - T b = 4; - res_access[0] = std::hypot(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceIlogb; - -template -void device_ilogb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::ilogb(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLgamma; - -template -void device_lgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = NAN; - res_access[0] = std::lgamma(a); - }); - }); - } - - assert(std::isnan(result)); -} - -template -class DeviceLog1p; - -template -void device_log1p_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0; - res_access[0] = std::log1p(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLog2; - -template -void device_log2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::log2(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceLogb; - -template -void device_logb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 1; - res_access[0] = std::logb(a); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceRemainder; - -template -void device_remainder_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; - T ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0.5; - T b = 1; - res_access[0] = std::remainder(a, b); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceRemquo; - -template -void device_remquo_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; int quo = -1; - T ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&quo, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto quo_access = buffer2.template get_access(cgh); - cgh.single_task >([=]() { - T a = 0.5; - T b = 1; - res_access[0] = std::remquo(a, b, &quo_access[0]); - }); - }); - } - - assert(result == ref); -} - -template -class DeviceTgamma; - -template -void device_tgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -1; { - s::buffer buffer1(&result, numOfItems); + s::buffer buffer1(result, numOfItems); + s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer3(&iptr, s::range<1>{1}); + s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - T a = NAN; - res_access[0] = std::tgamma(a); - }); - }); - } - - assert(std::isnan(result)); -} - -template -void device_cmath_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); - device_acos_test(deviceQueue); - device_asin_test(deviceQueue); - device_atan_test(deviceQueue); - device_atan2_test(deviceQueue); - device_cosh_test(deviceQueue); - device_exp_test(deviceQueue); - device_fmod_test(deviceQueue); - device_frexp_test(deviceQueue); - device_ldexp_test(deviceQueue); - device_log10_test(deviceQueue); - device_modf_test(deviceQueue); - device_pow_test(deviceQueue); - device_sinh_test(deviceQueue); - device_sqrt_test(deviceQueue); - device_tan_test(deviceQueue); - device_tanh_test(deviceQueue); - device_acosh_test(deviceQueue); - device_asinh_test(deviceQueue); - device_atanh_test(deviceQueue); - device_cbrt_test(deviceQueue); - device_erf_test(deviceQueue); - device_erfc_test(deviceQueue); - device_exp2_test(deviceQueue); - device_expm1_test(deviceQueue); - device_fdim_test(deviceQueue); - device_fma_test(deviceQueue); - device_hypot_test(deviceQueue); - device_ilogb_test(deviceQueue); - device_lgamma_test(deviceQueue); - device_log1p_test(deviceQueue); - device_log2_test(deviceQueue); - device_logb_test(deviceQueue); - device_remainder_test(deviceQueue); - device_remquo_test(deviceQueue); - device_tgamma_test(deviceQueue); + auto exp_access = buffer2.template get_access(cgh); + auto iptr_access = buffer3.template get_access(cgh); + auto quo_access = buffer4.template get_access(cgh); + cgh.single_task([=]() { + int i = 0; + { + T a = 0; + res_access[i++] = std::cos(a); + } + { + T a = 0; + res_access[i++] = std::sin(a); + } + { + T a = 1; + res_access[i++] = std::log(a); + } + { + T a = 1; + res_access[i++] = std::acos(a); + } + { + T a = 0; + res_access[i++] = std::asin(a); + } + { + T a = 0; + res_access[i++] = std::atan(a); + } + { + T a = 0; + T b = 1; + res_access[i++] = std::atan2(a, b); + } + { + T a = 0; + res_access[i++] = std::cosh(a); + } + { + T a = 0; + res_access[i++] = std::exp(a); + } + { + T a = 1.5; + T b = 1; + res_access[i++] = std::fmod(a, b); + } + { + T a = 0; + res_access[i++] = std::frexp(a, &exp_access[0]); + } + { + T a = 1; + res_access[i++] = std::ldexp(a, 1); + } + { + T a = 1; + res_access[i++] = std::log10(a); + } + { + T a = 1; + res_access[i++] = std::modf(a, &iptr_access[0]); + } + { + T a = 1; + T b = 1; + res_access[i++] = std::pow(a, b); + } + { + T a = 0; + res_access[i++] = std::sinh(a); + } + { + T a = 4; + res_access[i++] = std::sqrt(a); + } + { + T a = 0; + res_access[i++] = std::tan(a); + } + { + T a = 0; + res_access[i++] = std::tanh(a); + } + { + T a = 1; + res_access[i++] = std::acosh(a); + } + { + T a = 0; + res_access[i++] = std::asinh(a); + } + { + T a = 0; + res_access[i++] = std::atanh(a); + } + { + T a = 1; + res_access[i++] = std::cbrt(a); + } + { + T a = 0; + res_access[i++] = std::erf(a); + } + { + T a = 0; + res_access[i++] = std::erfc(a); + } + { + T a = 1; + res_access[i++] = std::exp2(a); + } + { + T a = 0; + res_access[i++] = std::expm1(a); + } + { + T a = 0; + res_access[i++] = std::fdim(1, a); + } + { + T a = 1; + T b = 1; + T c = 1; + res_access[i++] = std::fma(a, b, c); + } + { + T a = 3; + T b = 4; + res_access[i++] = std::hypot(a, b); + } + { + T a = 1; + res_access[i++] = std::ilogb(a); + } + { + T a = 0; + res_access[i++] = std::log1p(a); + } + { + T a = 1; + res_access[i++] = std::log2(a); + } + { + T a = 1; + res_access[i++] = std::logb(a); + } + { + T a = 0.5; + T b = 1; + res_access[i++] = std::remainder(a, b); + } + { + T a = 0.5; + T b = 1; + res_access[i++] = std::remquo(a, b, &quo_access[0]); + } + { + T a = NAN; + res_access[i++] = std::tgamma(a); + } + { + T a = NAN; + res_access[i++] = std::lgamma(a); + } + }); + }); + } + for (int i = 0; i < 36; ++i) { + assert(is_about_FP(result[i], ref[i])); + } + assert(std::isnan(result[36])); + assert(std::isnan(result[37])); + assert(is_about_FP(iptr, b)); + assert(expv == 0); + assert(quo == 0); } int main() { diff --git a/sycl/test/devicelib/math_fp64_test.cpp b/sycl/test/devicelib/math_fp64_test.cpp index a56da6dea2301..274055ce46419 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -4,825 +4,247 @@ #include #include #include +#include "math_utils.hpp" namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -class DeviceCos; +double a = 0; +double b = 1; +double c = 0.5; +double d = 2; +double e = 5; -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = cos(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceSin; - -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = sin(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLog; - -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = log(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAcos; - -void device_acos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = acos(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAsin; - -void device_asin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = asin(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAtan; - -void device_atan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = atan(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAtan2; - -void device_atan2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - double b = 1; - res_access[0] = atan2(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceCosh; - -void device_cosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = cosh(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceExp; - -void device_exp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = exp(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceFmod; - -void device_fmod_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1.5; - double b = 1; - res_access[0] = fmod(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceFrexp; - -void device_frexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - int exp = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&exp, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto exp_access = buffer2.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = frexp(a, &exp_access[0]); - }); - }); - } - - assert(result == ref && exp == 0); -} - -class DeviceLdexp; - -void device_ldexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = ldexp(a, 1); - }); - }); - } - - assert(result == ref); -} - -class DeviceLog10; - -void device_log10_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = log10(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceModf; - -void device_modf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; +void device_math_test(s::queue &deviceQueue) { + s::range<1> numOfItems{38}; + double result[38] = {-1}; + double ref[38] = { + b, + a, + a, + a, + a, + a, + a, + b, + b, + c, + a, + d, + a, + a, + b, + a, + d, + a, + a, + a, + a, + a, + b, + a, + b, + d, + a, + b, + d, + e, + a, + a, + a, + a, + c, + c, + a, + a, + }; + int expv = -1; double iptr = -1; - double ref1 = 0; - double ref2 = 1; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&iptr, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto iptr_access = buffer2.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = modf(a, &iptr_access[0]); - }); - }); - } - - assert(result == ref1 && iptr == ref2); -} - -class DevicePow; - -void device_pow_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - double b = 1; - res_access[0] = pow(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceSinh; - -void device_sinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = sinh(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceSqrt; - -void device_sqrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 4; - res_access[0] = sqrt(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceTan; - -void device_tan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = tan(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceTanh; - -void device_tanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = tanh(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAcosh; - -void device_acosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = acosh(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAsinh; - -void device_asinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = asinh(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAtanh; - -void device_atanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = atanh(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceCbrt; - -void device_cbrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = cbrt(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceErf; - -void device_erf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = erf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceErfc; - -void device_erfc_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = erfc(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceExp2; - -void device_exp2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = exp2(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceExpm1; - -void device_expm1_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = expm1(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceFdim; - -void device_fdim_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - double b = 0; - res_access[0] = fdim(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceFma; - -void device_fma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - double b = 1; - double c = 1; - res_access[0] = fma(a, b, c); - }); - }); - } - - assert(result == ref); -} - -class DeviceHypot; - -void device_hypot_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 3; - double b = 4; - res_access[0] = hypot(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceIlogb; - -void device_ilogb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = ilogb(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLgamma; - -void device_lgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = NAN; - res_access[0] = lgamma(a); - }); - }); - } - - assert(std::isnan(result)); -} - -class DeviceLog1p; - -void device_log1p_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0; - res_access[0] = log1p(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLog2; - -void device_log2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = log2(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLogb; - -void device_logb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 1; - res_access[0] = logb(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceRemainder; - -void device_remainder_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - double ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = 0.5; - double b = 1; - res_access[0] = remainder(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceRemquo; - -void device_remquo_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; int quo = -1; - double ref = 0.5; { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&quo, numOfItems); + s::buffer buffer1(result, numOfItems); + s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer3(&iptr, s::range<1>{1}); + s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); - auto quo_access = buffer2.template get_access(cgh); - cgh.single_task([=]() { - double a = 0.5; - double b = 1; - res_access[0] = remquo(a, b, &quo_access[0]); - }); - }); - } - - assert(result == ref); -} - -class DeviceTgamma; -void device_tgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - double a = NAN; - res_access[0] = tgamma(a); - }); - }); - } - - assert(std::isnan(result)); -} -void device_math_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); - device_acos_test(deviceQueue); - device_asin_test(deviceQueue); - device_atan_test(deviceQueue); - device_atan2_test(deviceQueue); - device_cosh_test(deviceQueue); - device_exp_test(deviceQueue); - device_fmod_test(deviceQueue); - device_frexp_test(deviceQueue); - device_ldexp_test(deviceQueue); - device_log10_test(deviceQueue); - device_modf_test(deviceQueue); - device_pow_test(deviceQueue); - device_sinh_test(deviceQueue); - device_sqrt_test(deviceQueue); - device_tan_test(deviceQueue); - device_tanh_test(deviceQueue); - device_acosh_test(deviceQueue); - device_asinh_test(deviceQueue); - device_atanh_test(deviceQueue); - device_cbrt_test(deviceQueue); - device_erf_test(deviceQueue); - device_erfc_test(deviceQueue); - device_exp2_test(deviceQueue); - device_expm1_test(deviceQueue); - device_fdim_test(deviceQueue); - device_fma_test(deviceQueue); - device_hypot_test(deviceQueue); - device_ilogb_test(deviceQueue); - device_lgamma_test(deviceQueue); - device_log1p_test(deviceQueue); - device_log2_test(deviceQueue); - device_logb_test(deviceQueue); - device_remainder_test(deviceQueue); - device_remquo_test(deviceQueue); - device_tgamma_test(deviceQueue); + auto exp_access = buffer2.template get_access(cgh); + auto iptr_access = buffer3.template get_access(cgh); + auto quo_access = buffer4.template get_access(cgh); + cgh.single_task([=]() { + int i = 0; + { + double a = 0; + res_access[i++] = cos(a); + } + { + double a = 0; + res_access[i++] = sin(a); + } + { + double a = 1; + res_access[i++] = log(a); + } + { + double a = 1; + res_access[i++] = acos(a); + } + { + double a = 0; + res_access[i++] = asin(a); + } + { + double a = 0; + res_access[i++] = atan(a); + } + { + double a = 0; + double b = 1; + res_access[i++] = atan2(a, b); + } + { + double a = 0; + res_access[i++] = cosh(a); + } + { + double a = 0; + res_access[i++] = exp(a); + } + { + double a = 1.5; + double b = 1; + res_access[i++] = fmod(a, b); + } + { + double a = 0; + res_access[i++] = frexp(a, &exp_access[0]); + } + { + double a = 1; + res_access[i++] = ldexp(a, 1); + } + { + double a = 1; + res_access[i++] = log10(a); + } + { + double a = 1; + res_access[i++] = modf(a, &iptr_access[0]); + } + { + double a = 1; + double b = 1; + res_access[i++] = pow(a, b); + } + { + double a = 0; + res_access[i++] = sinh(a); + } + { + double a = 4; + res_access[i++] = sqrt(a); + } + { + double a = 0; + res_access[i++] = tan(a); + } + { + double a = 0; + res_access[i++] = tanh(a); + } + { + double a = 1; + res_access[i++] = acosh(a); + } + { + double a = 0; + res_access[i++] = asinh(a); + } + { + double a = 0; + res_access[i++] = atanh(a); + } + { + double a = 1; + res_access[i++] = cbrt(a); + } + { + double a = 0; + res_access[i++] = erf(a); + } + { + double a = 0; + res_access[i++] = erfc(a); + } + { + double a = 1; + res_access[i++] = exp2(a); + } + { + double a = 0; + res_access[i++] = expm1(a); + } + { + double a = 0; + res_access[i++] = fdim(1, a); + } + { + double a = 1; + double b = 1; + double c = 1; + res_access[i++] = fma(a, b, c); + } + { + double a = 3; + double b = 4; + res_access[i++] = hypot(a, b); + } + { + double a = 1; + res_access[i++] = ilogb(a); + } + { + double a = 0; + res_access[i++] = log1p(a); + } + { + double a = 1; + res_access[i++] = log2(a); + } + { + double a = 1; + res_access[i++] = logb(a); + } + { + double a = 0.5; + double b = 1; + res_access[i++] = remainder(a, b); + } + { + double a = 0.5; + double b = 1; + res_access[i++] = remquo(a, b, &quo_access[0]); + } + { + double a = NAN; + res_access[i++] = tgamma(a); + } + { + double a = NAN; + res_access[i++] = lgamma(a); + } + }); + }); + } + for (int i = 0; i < 36; ++i) { + assert(is_about_FP(result[i], ref[i])); + } + assert(std::isnan(result[36])); + assert(std::isnan(result[37])); + assert(is_about_FP(iptr, b)); + assert(expv == 0); + assert(quo == 0); } int main() { diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index f60f376b52443..90589040f4b70 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -4,826 +4,247 @@ #include #include #include +#include "math_utils.hpp" namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -class DeviceCos; +float a = 0; +float b = 1; +float c = 0.5; +float d = 2; +float e = 5; -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = cosf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceSin; - -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = sinf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLog; - -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = logf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAcos; - -void device_acos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = acosf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAsin; - -void device_asin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = asinf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAtan; - -void device_atan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = atanf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAtan2; - -void device_atan2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - float b = 1; - res_access[0] = atan2f(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceCosh; - -void device_cosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = coshf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceExp; - -void device_exp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = expf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceFmod; - -void device_fmod_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1.5; - float b = 1; - res_access[0] = fmodf(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceFrexp; - -void device_frexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - int exp = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&exp, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto exp_access = buffer2.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = frexpf(a, &exp_access[0]); - }); - }); - } - - assert(result == ref && exp == 0); -} - -class DeviceLdexp; - -void device_ldexp_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = ldexpf(a, 1); - }); - }); - } - - assert(result == ref); -} - -class DeviceLog10; - -void device_log10_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = log10f(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceModf; - -void device_modf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; +void device_math_test(s::queue &deviceQueue) { + s::range<1> numOfItems{38}; + float result[38] = {-1}; + float ref[38] = { + b, + a, + a, + a, + a, + a, + a, + b, + b, + c, + a, + d, + a, + a, + b, + a, + d, + a, + a, + a, + a, + a, + b, + a, + b, + d, + a, + b, + d, + e, + a, + a, + a, + a, + c, + c, + a, + a, + }; + int expv = -1; float iptr = -1; - float ref1 = 0; - float ref2 = 1; - { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&iptr, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - auto iptr_access = buffer2.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = modff(a, &iptr_access[0]); - }); - }); - } - - assert(result == ref1 && iptr == ref2); -} - -class DevicePow; - -void device_pow_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - float b = 1; - res_access[0] = powf(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceSinh; - -void device_sinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = sinhf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceSqrt; - -void device_sqrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 4; - res_access[0] = sqrtf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceTan; - -void device_tan_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = tanf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceTanh; - -void device_tanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = tanhf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAcosh; - -void device_acosh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = acoshf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAsinh; - -void device_asinh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = asinhf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceAtanh; - -void device_atanh_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = atanhf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceCbrt; - -void device_cbrt_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = cbrtf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceErf; - -void device_erf_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = erff(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceErfc; - -void device_erfc_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = erfcf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceExp2; - -void device_exp2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = exp2f(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceExpm1; - -void device_expm1_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = expm1f(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceFdim; - -void device_fdim_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = fdimf(1, a); - }); - }); - } - - assert(result == ref); -} - -class DeviceFma; - -void device_fma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 2; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - float b = 1; - float c = 1; - res_access[0] = fmaf(a, b, c); - }); - }); - } - - assert(result == ref); -} - -class DeviceHypot; - -void device_hypot_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 3; - float b = 4; - res_access[0] = hypotf(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceIlogb; - -void device_ilogb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = ilogbf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLgamma; - -void device_lgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = NAN; - res_access[0] = lgammaf(a); - }); - }); - } - - assert(std::isnan(result)); -} - -class DeviceLog1p; - -void device_log1p_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0; - res_access[0] = log1pf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLog2; - -void device_log2_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = log2f(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceLogb; - -void device_logb_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 1; - res_access[0] = logbf(a); - }); - }); - } - - assert(result == ref); -} - -class DeviceRemainder; - -void device_remainder_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - float ref = 0.5; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = 0.5; - float b = 1; - res_access[0] = remainderf(a, b); - }); - }); - } - - assert(result == ref); -} - -class DeviceRemquo; - -void device_remquo_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; int quo = -1; - float ref = 0.5; { - s::buffer buffer1(&result, numOfItems); - s::buffer buffer2(&quo, numOfItems); + s::buffer buffer1(result, numOfItems); + s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer3(&iptr, s::range<1>{1}); + s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { auto res_access = buffer1.template get_access(cgh); - auto quo_access = buffer2.template get_access(cgh); - cgh.single_task([=]() { - float a = 0.5; - float b = 1; - res_access[0] = remquof(a, b, &quo_access[0]); - }); - }); - } - - assert(result == ref); -} - -class DeviceTgamma; - -void device_tgamma_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result = -1; - { - s::buffer buffer1(&result, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access = buffer1.template get_access(cgh); - cgh.single_task([=]() { - float a = NAN; - res_access[0] = tgammaf(a); - }); - }); - } - - assert(std::isnan(result)); -} - -void device_math_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); - device_acos_test(deviceQueue); - device_asin_test(deviceQueue); - device_atan_test(deviceQueue); - device_atan2_test(deviceQueue); - device_cosh_test(deviceQueue); - device_exp_test(deviceQueue); - device_fmod_test(deviceQueue); - device_frexp_test(deviceQueue); - device_ldexp_test(deviceQueue); - device_log10_test(deviceQueue); - device_modf_test(deviceQueue); - device_pow_test(deviceQueue); - device_sinh_test(deviceQueue); - device_sqrt_test(deviceQueue); - device_tan_test(deviceQueue); - device_tanh_test(deviceQueue); - device_acosh_test(deviceQueue); - device_asinh_test(deviceQueue); - device_atanh_test(deviceQueue); - device_cbrt_test(deviceQueue); - device_erf_test(deviceQueue); - device_erfc_test(deviceQueue); - device_exp2_test(deviceQueue); - device_expm1_test(deviceQueue); - device_fdim_test(deviceQueue); - device_fma_test(deviceQueue); - device_hypot_test(deviceQueue); - device_ilogb_test(deviceQueue); - device_lgamma_test(deviceQueue); - device_log1p_test(deviceQueue); - device_log2_test(deviceQueue); - device_logb_test(deviceQueue); - device_remainder_test(deviceQueue); - device_remquo_test(deviceQueue); - device_tgamma_test(deviceQueue); + auto exp_access = buffer2.template get_access(cgh); + auto iptr_access = buffer3.template get_access(cgh); + auto quo_access = buffer4.template get_access(cgh); + cgh.single_task([=]() { + int i = 0; + { + float a = 0; + res_access[i++] = cosf(a); + } + { + float a = 0; + res_access[i++] = sinf(a); + } + { + float a = 1; + res_access[i++] = logf(a); + } + { + float a = 1; + res_access[i++] = acosf(a); + } + { + float a = 0; + res_access[i++] = asinf(a); + } + { + float a = 0; + res_access[i++] = atanf(a); + } + { + float a = 0; + float b = 1; + res_access[i++] = atan2f(a, b); + } + { + float a = 0; + res_access[i++] = coshf(a); + } + { + float a = 0; + res_access[i++] = expf(a); + } + { + float a = 1.5; + float b = 1; + res_access[i++] = fmodf(a, b); + } + { + float a = 0; + res_access[i++] = frexpf(a, &exp_access[0]); + } + { + float a = 1; + res_access[i++] = ldexpf(a, 1); + } + { + float a = 1; + res_access[i++] = log10f(a); + } + { + float a = 1; + res_access[i++] = modff(a, &iptr_access[0]); + } + { + float a = 1; + float b = 1; + res_access[i++] = powf(a, b); + } + { + float a = 0; + res_access[i++] = sinhf(a); + } + { + float a = 4; + res_access[i++] = sqrtf(a); + } + { + float a = 0; + res_access[i++] = tanf(a); + } + { + float a = 0; + res_access[i++] = tanhf(a); + } + { + float a = 1; + res_access[i++] = acoshf(a); + } + { + float a = 0; + res_access[i++] = asinhf(a); + } + { + float a = 0; + res_access[i++] = atanhf(a); + } + { + float a = 1; + res_access[i++] = cbrtf(a); + } + { + float a = 0; + res_access[i++] = erff(a); + } + { + float a = 0; + res_access[i++] = erfcf(a); + } + { + float a = 1; + res_access[i++] = exp2f(a); + } + { + float a = 0; + res_access[i++] = expm1f(a); + } + { + float a = 0; + res_access[i++] = fdimf(1, a); + } + { + float a = 1; + float b = 1; + float c = 1; + res_access[i++] = fmaf(a, b, c); + } + { + float a = 3; + float b = 4; + res_access[i++] = hypotf(a, b); + } + { + float a = 1; + res_access[i++] = ilogbf(a); + } + { + float a = 0; + res_access[i++] = log1pf(a); + } + { + float a = 1; + res_access[i++] = log2f(a); + } + { + float a = 1; + res_access[i++] = logbf(a); + } + { + float a = 0.5; + float b = 1; + res_access[i++] = remainderf(a, b); + } + { + float a = 0.5; + float b = 1; + res_access[i++] = remquof(a, b, &quo_access[0]); + } + { + float a = NAN; + res_access[i++] = tgammaf(a); + } + { + float a = NAN; + res_access[i++] = lgammaf(a); + } + }); + }); + } + for (int i = 0; i < 36; ++i) { + assert(is_about_FP(result[i], ref[i])); + } + assert(std::isnan(result[36])); + assert(std::isnan(result[37])); + assert(is_about_FP(iptr, b)); + assert(expv == 0); + assert(quo == 0); } int main() { From 79dce3e56fe1752b4e7b409a60f072c68bee0696 Mon Sep 17 00:00:00 2001 From: haonanya Date: Mon, 9 Mar 2020 20:45:19 +0800 Subject: [PATCH 04/19] [SYCL] Remove using namespace std in math_utils.hpp Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index ef13d36a54ba9..e15837bfbf920 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -1,7 +1,6 @@ #ifndef MATH_UTILS #include #include -using namespace std; // T must be float-point type template bool is_about_FP(T x, T y) { @@ -9,15 +8,15 @@ bool is_about_FP(T x, T y) { return true; else { if (x != 0 && y != 0) { - T max_v = fmax(abs(x), abs(y)); - return (abs(x - y) / max_v) < - numeric_limits::epsilon() * 100; + T max_v = std::fmax(std::abs(x), std::abs(y)); + return (std::abs(x - y) / max_v) < + std::numeric_limits::epsilon() * 100; } else { if (x != 0) - return abs(x) < numeric_limits::epsilon() * 100; + return std::abs(x) < std::numeric_limits::epsilon() * 100; else - return abs(y) < numeric_limits::epsilon() * 100; + return std::abs(y) < std::numeric_limits::epsilon() * 100; } } } From 3836b41a4400925a20ac0d07ae30741a02d6b9da Mon Sep 17 00:00:00 2001 From: haonanya Date: Mon, 9 Mar 2020 20:47:19 +0800 Subject: [PATCH 05/19] [SYCL] modify std_complex_math_fp64_test.cpp and std_complex_math_test.cpp, as math_utils.hpp changes Signed-off-by: haonanya --- sycl/test/devicelib/std_complex_math_fp64_test.cpp | 1 + sycl/test/devicelib/std_complex_math_test.cpp | 2 ++ 2 files changed, 3 insertions(+) diff --git a/sycl/test/devicelib/std_complex_math_fp64_test.cpp b/sycl/test/devicelib/std_complex_math_fp64_test.cpp index e1213a713ad5b..31e0f65a388fc 100644 --- a/sycl/test/devicelib/std_complex_math_fp64_test.cpp +++ b/sycl/test/devicelib/std_complex_math_fp64_test.cpp @@ -4,6 +4,7 @@ #include #include #include "math_utils.hpp" +using namespace std; namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; diff --git a/sycl/test/devicelib/std_complex_math_test.cpp b/sycl/test/devicelib/std_complex_math_test.cpp index f3fe5ae9ae510..af4ef799fa2f4 100644 --- a/sycl/test/devicelib/std_complex_math_test.cpp +++ b/sycl/test/devicelib/std_complex_math_test.cpp @@ -5,6 +5,8 @@ #include #include "math_utils.hpp" +using namespace std; + namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; From e28efb4c005f5a36663d473141de8eaa2e0a8e9b Mon Sep 17 00:00:00 2001 From: haonanya Date: Wed, 11 Mar 2020 10:32:10 +0800 Subject: [PATCH 06/19] [SYCL] [Test] Put all math function calls into a single kernel Signed-off-by: haonanya --- sycl/test/devicelib/cmath_fp64_test.cpp | 82 ++++++++----------------- sycl/test/devicelib/cmath_test.cpp | 82 ++++++++----------------- sycl/test/devicelib/math_fp64_test.cpp | 82 ++++++++----------------- sycl/test/devicelib/math_test.cpp | 82 ++++++++----------------- 4 files changed, 104 insertions(+), 224 deletions(-) diff --git a/sycl/test/devicelib/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp index bbe6a39cf819b..b28fb3ba67a9c 100644 --- a/sycl/test/devicelib/cmath_fp64_test.cpp +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -10,62 +10,29 @@ namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -double a = 0; -double b = 1; -double c = 0.5; -double d = 2; -double e = 5; +#define TEST_NUM 38 + +double ref[TEST_NUM] = { +1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5, +0, 2, 0, 0, 1, 0, 2, 0, 0, 0, +0, 0, 1, 0, 1, 2, 0, 1, 2, 5, +0, 0, 0, 0, 0.5, 0.5, NAN, NAN,}; + +double refIptr = 1; template void device_cmath_test(s::queue &deviceQueue) { - s::range<1> numOfItems{38}; - T result[38] = {-1}; - T ref[38] = { - b, - a, - a, - a, - a, - a, - a, - b, - b, - c, - a, - d, - a, - a, - b, - a, - d, - a, - a, - a, - a, - a, - b, - a, - b, - d, - a, - b, - d, - e, - a, - a, - a, - a, - c, - c, - a, - a, - }; - int expv = -1; + s::range<1> numOfItems{TEST_NUM}; + T result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; + // Variable iptr stores the integral part of float point in modf function T iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { s::buffer buffer1(result, numOfItems); - s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer2(&exponent, s::range<1>{1}); s::buffer buffer3(&iptr, s::range<1>{1}); s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { @@ -187,8 +154,9 @@ void device_cmath_test(s::queue &deviceQueue) { res_access[i++] = std::expm1(a); } { - T a = 0; - res_access[i++] = std::fdim(1, a); + T a = 1; + T b = 0; + res_access[i++] = std::fdim(a, b); } { T a = 1; @@ -238,13 +206,15 @@ void device_cmath_test(s::queue &deviceQueue) { }); }); } - for (int i = 0; i < 36; ++i) { + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } - assert(std::isnan(result[36])); - assert(std::isnan(result[37])); - assert(is_about_FP(iptr, b)); - assert(expv == 0); + // Test modf integral part + assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent + assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/cmath_test.cpp b/sycl/test/devicelib/cmath_test.cpp index 7f6866e0a8b61..fa96ff238e10d 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -10,62 +10,29 @@ namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -float a = 0; -float b = 1; -float c = 0.5; -float d = 2; -float e = 5; +#define TEST_NUM 38 + +float ref[TEST_NUM] = { +1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5, +0, 2, 0, 0, 1, 0, 2, 0, 0, 0, +0, 0, 1, 0, 1, 2, 0, 1, 2, 5, +0, 0, 0, 0, 0.5, 0.5, NAN, NAN,}; + +float refIptr = 1; template void device_cmath_test(s::queue &deviceQueue) { - s::range<1> numOfItems{38}; - T result[38] = {-1}; - T ref[38] = { - b, - a, - a, - a, - a, - a, - a, - b, - b, - c, - a, - d, - a, - a, - b, - a, - d, - a, - a, - a, - a, - a, - b, - a, - b, - d, - a, - b, - d, - e, - a, - a, - a, - a, - c, - c, - a, - a, - }; - int expv = -1; + s::range<1> numOfItems{TEST_NUM}; + T result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; + // Variable iptr stores the integral part of float point in modf function T iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { s::buffer buffer1(result, numOfItems); - s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer2(&exponent, s::range<1>{1}); s::buffer buffer3(&iptr, s::range<1>{1}); s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { @@ -187,8 +154,9 @@ void device_cmath_test(s::queue &deviceQueue) { res_access[i++] = std::expm1(a); } { - T a = 0; - res_access[i++] = std::fdim(1, a); + T a = 1; + T b = 0; + res_access[i++] = std::fdim(a, b); } { T a = 1; @@ -238,13 +206,15 @@ void device_cmath_test(s::queue &deviceQueue) { }); }); } - for (int i = 0; i < 36; ++i) { + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } - assert(std::isnan(result[36])); - assert(std::isnan(result[37])); - assert(is_about_FP(iptr, b)); - assert(expv == 0); + // Test modf integral part + assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent + assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/math_fp64_test.cpp b/sycl/test/devicelib/math_fp64_test.cpp index 274055ce46419..5b9eb95695b20 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -10,61 +10,28 @@ namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -double a = 0; -double b = 1; -double c = 0.5; -double d = 2; -double e = 5; +#define TEST_NUM 38 + +double ref[TEST_NUM] = { +1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5, +0, 2, 0, 0, 1, 0, 2, 0, 0, 0, +0, 0, 1, 0, 1, 2, 0, 1, 2, 5, +0, 0, 0, 0, 0.5, 0.5, NAN, NAN,}; + +double refIptr = 1; void device_math_test(s::queue &deviceQueue) { - s::range<1> numOfItems{38}; - double result[38] = {-1}; - double ref[38] = { - b, - a, - a, - a, - a, - a, - a, - b, - b, - c, - a, - d, - a, - a, - b, - a, - d, - a, - a, - a, - a, - a, - b, - a, - b, - d, - a, - b, - d, - e, - a, - a, - a, - a, - c, - c, - a, - a, - }; - int expv = -1; + s::range<1> numOfItems{TEST_NUM}; + double result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; + // Variable iptr stores the integral part of float point in modf function double iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { s::buffer buffer1(result, numOfItems); - s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer2(&exponent, s::range<1>{1}); s::buffer buffer3(&iptr, s::range<1>{1}); s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { @@ -186,8 +153,9 @@ void device_math_test(s::queue &deviceQueue) { res_access[i++] = expm1(a); } { - double a = 0; - res_access[i++] = fdim(1, a); + double a = 1; + double b = 0; + res_access[i++] = fdim(a, b); } { double a = 1; @@ -237,13 +205,15 @@ void device_math_test(s::queue &deviceQueue) { }); }); } - for (int i = 0; i < 36; ++i) { + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } - assert(std::isnan(result[36])); - assert(std::isnan(result[37])); - assert(is_about_FP(iptr, b)); - assert(expv == 0); + // Test modf integral part + assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent + assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index 90589040f4b70..6ec771c941a09 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -10,61 +10,28 @@ namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; -float a = 0; -float b = 1; -float c = 0.5; -float d = 2; -float e = 5; +#define TEST_NUM 38 + +float ref[TEST_NUM] = { +1, 0, 0, 0, 0, 0, 0, 1, 1, 0.5, +0, 2, 0, 0, 1, 0, 2, 0, 0, 0, +0, 0, 1, 0, 1, 2, 0, 1, 2, 5, +0, 0, 0, 0, 0.5, 0.5, NAN, NAN,}; + +float refIptr = 1; void device_math_test(s::queue &deviceQueue) { - s::range<1> numOfItems{38}; - float result[38] = {-1}; - float ref[38] = { - b, - a, - a, - a, - a, - a, - a, - b, - b, - c, - a, - d, - a, - a, - b, - a, - d, - a, - a, - a, - a, - a, - b, - a, - b, - d, - a, - b, - d, - e, - a, - a, - a, - a, - c, - c, - a, - a, - }; - int expv = -1; + s::range<1> numOfItems{TEST_NUM}; + float result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; + // Variable iptr stores the integral part of float point in modf function float iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { s::buffer buffer1(result, numOfItems); - s::buffer buffer2(&expv, s::range<1>{1}); + s::buffer buffer2(&exponent, s::range<1>{1}); s::buffer buffer3(&iptr, s::range<1>{1}); s::buffer buffer4(&quo, s::range<1>{1}); deviceQueue.submit([&](cl::sycl::handler &cgh) { @@ -186,8 +153,9 @@ void device_math_test(s::queue &deviceQueue) { res_access[i++] = expm1f(a); } { - float a = 0; - res_access[i++] = fdimf(1, a); + float a = 1; + float b = 0; + res_access[i++] = fdimf(a, b); } { float a = 1; @@ -237,13 +205,15 @@ void device_math_test(s::queue &deviceQueue) { }); }); } - for (int i = 0; i < 36; ++i) { + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } - assert(std::isnan(result[36])); - assert(std::isnan(result[37])); - assert(is_about_FP(iptr, b)); - assert(expv == 0); + // Test modf integral part + assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent + assert(exponent == 0); + // Test remquo sign assert(quo == 0); } From 1251e565f7c08d677af2a78f52f839e1a09acbdc Mon Sep 17 00:00:00 2001 From: haonanya Date: Wed, 11 Mar 2020 10:37:28 +0800 Subject: [PATCH 07/19] [SYCL] [Test] Deal with nan and inf input in is_about_FP function in math_utils.hpp Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index e15837bfbf920..22c0c17b4e8a0 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -4,15 +4,19 @@ // T must be float-point type template bool is_about_FP(T x, T y) { - if (x == y) - return true; + // At least one input is nan + if (std::isnan(x) || std::isnan(y)) + return std::isnan(x) && std::isnan(y); + // At least one input is inf + else if (std::isinf(x) || std::isinf(y)) + return (x == y); + // two finite else { if (x != 0 && y != 0) { T max_v = std::fmax(std::abs(x), std::abs(y)); return (std::abs(x - y) / max_v) < - std::numeric_limits::epsilon() * 100; - } - else { + std::numeric_limits::epsilon() * 100; + } else { if (x != 0) return std::abs(x) < std::numeric_limits::epsilon() * 100; else From 9d62247fbfac4cfa310b1e8fb397d6719566385e Mon Sep 17 00:00:00 2001 From: haonanya Date: Wed, 11 Mar 2020 13:27:20 +0800 Subject: [PATCH 08/19] [SYCL] Format code Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index 22c0c17b4e8a0..d2de1cb508911 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -7,10 +7,12 @@ bool is_about_FP(T x, T y) { // At least one input is nan if (std::isnan(x) || std::isnan(y)) return std::isnan(x) && std::isnan(y); - // At least one input is inf + + // At least one input is inf else if (std::isinf(x) || std::isinf(y)) return (x == y); - // two finite + + // two finite else { if (x != 0 && y != 0) { T max_v = std::fmax(std::abs(x), std::abs(y)); From b6700560ff3fe6a7857c02eaa88e2aa2518d8d92 Mon Sep 17 00:00:00 2001 From: haonanya Date: Fri, 13 Mar 2020 15:33:59 +0800 Subject: [PATCH 09/19] [SYCL] Code clean Signed-off-by: haonanya --- sycl/test/devicelib/cmath_fp64_test.cpp | 207 ++++++----------------- sycl/test/devicelib/cmath_test.cpp | 207 ++++++----------------- sycl/test/devicelib/math_fp64_test.cpp | 209 ++++++------------------ sycl/test/devicelib/math_test.cpp | 207 ++++++----------------- sycl/test/devicelib/math_utils.hpp | 15 +- 5 files changed, 194 insertions(+), 651 deletions(-) diff --git a/sycl/test/devicelib/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp index b28fb3ba67a9c..294660ca8fc46 100644 --- a/sycl/test/devicelib/cmath_fp64_test.cpp +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -24,10 +24,13 @@ template void device_cmath_test(s::queue &deviceQueue) { s::range<1> numOfItems{TEST_NUM}; T result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function int exponent = -1; + // Variable iptr stores the integral part of float point in modf function T iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { @@ -42,178 +45,60 @@ void device_cmath_test(s::queue &deviceQueue) { auto quo_access = buffer4.template get_access(cgh); cgh.single_task([=]() { int i = 0; - { - T a = 0; - res_access[i++] = std::cos(a); - } - { - T a = 0; - res_access[i++] = std::sin(a); - } - { - T a = 1; - res_access[i++] = std::log(a); - } - { - T a = 1; - res_access[i++] = std::acos(a); - } - { - T a = 0; - res_access[i++] = std::asin(a); - } - { - T a = 0; - res_access[i++] = std::atan(a); - } - { - T a = 0; - T b = 1; - res_access[i++] = std::atan2(a, b); - } - { - T a = 0; - res_access[i++] = std::cosh(a); - } - { - T a = 0; - res_access[i++] = std::exp(a); - } - { - T a = 1.5; - T b = 1; - res_access[i++] = std::fmod(a, b); - } - { - T a = 0; - res_access[i++] = std::frexp(a, &exp_access[0]); - } - { - T a = 1; - res_access[i++] = std::ldexp(a, 1); - } - { - T a = 1; - res_access[i++] = std::log10(a); - } - { - T a = 1; - res_access[i++] = std::modf(a, &iptr_access[0]); - } - { - T a = 1; - T b = 1; - res_access[i++] = std::pow(a, b); - } - { - T a = 0; - res_access[i++] = std::sinh(a); - } - { - T a = 4; - res_access[i++] = std::sqrt(a); - } - { - T a = 0; - res_access[i++] = std::tan(a); - } - { - T a = 0; - res_access[i++] = std::tanh(a); - } - { - T a = 1; - res_access[i++] = std::acosh(a); - } - { - T a = 0; - res_access[i++] = std::asinh(a); - } - { - T a = 0; - res_access[i++] = std::atanh(a); - } - { - T a = 1; - res_access[i++] = std::cbrt(a); - } - { - T a = 0; - res_access[i++] = std::erf(a); - } - { - T a = 0; - res_access[i++] = std::erfc(a); - } - { - T a = 1; - res_access[i++] = std::exp2(a); - } - { - T a = 0; - res_access[i++] = std::expm1(a); - } - { - T a = 1; - T b = 0; - res_access[i++] = std::fdim(a, b); - } - { - T a = 1; - T b = 1; - T c = 1; - res_access[i++] = std::fma(a, b, c); - } - { - T a = 3; - T b = 4; - res_access[i++] = std::hypot(a, b); - } - { - T a = 1; - res_access[i++] = std::ilogb(a); - } - { - T a = 0; - res_access[i++] = std::log1p(a); - } - { - T a = 1; - res_access[i++] = std::log2(a); - } - { - T a = 1; - res_access[i++] = std::logb(a); - } - { - T a = 0.5; - T b = 1; - res_access[i++] = std::remainder(a, b); - } - { - T a = 0.5; - T b = 1; - res_access[i++] = std::remquo(a, b, &quo_access[0]); - } - { - T a = NAN; - res_access[i++] = std::tgamma(a); - } - { - T a = NAN; - res_access[i++] = std::lgamma(a); - } + res_access[i++] = std::cos(0.0); + res_access[i++] = std::sin(0.0); + res_access[i++] = std::log(1.0); + res_access[i++] = std::acos(1.0); + res_access[i++] = std::asin(0.0); + res_access[i++] = std::atan(0.0); + res_access[i++] = std::atan2(0.0, 1.0); + res_access[i++] = std::cosh(0.0); + res_access[i++] = std::exp(0.0); + res_access[i++] = std::fmod(1.5, 1.0); + res_access[i++] = std::frexp(0.0, &exp_access[0]); + res_access[i++] = std::ldexp(1.0, 1); + res_access[i++] = std::log10(1.0); + res_access[i++] = std::modf(1.0, &iptr_access[0]); + res_access[i++] = std::pow(1.0, 1.0); + res_access[i++] = std::sinh(0.0); + res_access[i++] = std::sqrt(4.0); + res_access[i++] = std::tan(0.0); + res_access[i++] = std::tanh(0.0); + res_access[i++] = std::acosh(1.0); + res_access[i++] = std::asinh(0.0); + res_access[i++] = std::atanh(0.0); + res_access[i++] = std::cbrt(1.0); + res_access[i++] = std::erf(0.0); + res_access[i++] = std::erfc(0.0); + res_access[i++] = std::exp2(1.0); + res_access[i++] = std::expm1(0.0); + res_access[i++] = std::fdim(1.0, 0.0); + res_access[i++] = std::fma(1.0, 1.0, 1.0); + res_access[i++] = std::hypot(3.0, 4.0); + res_access[i++] = std::ilogb(1.0); + res_access[i++] = std::log1p(0.0); + res_access[i++] = std::log2(1.0); + res_access[i++] = std::logb(1.0); + res_access[i++] = std::remainder(0.5, 1.0); + res_access[i++] = std::remquo(0.5, 1.0, &quo_access[0]); + T a = NAN; + res_access[i++] = std::tgamma(a); + res_access[i++] = std::lgamma(a); }); }); } + // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } + // Test modf integral part assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/cmath_test.cpp b/sycl/test/devicelib/cmath_test.cpp index fa96ff238e10d..4624b4bccf346 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -24,10 +24,13 @@ template void device_cmath_test(s::queue &deviceQueue) { s::range<1> numOfItems{TEST_NUM}; T result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function int exponent = -1; + // Variable iptr stores the integral part of float point in modf function T iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { @@ -42,178 +45,60 @@ void device_cmath_test(s::queue &deviceQueue) { auto quo_access = buffer4.template get_access(cgh); cgh.single_task([=]() { int i = 0; - { - T a = 0; - res_access[i++] = std::cos(a); - } - { - T a = 0; - res_access[i++] = std::sin(a); - } - { - T a = 1; - res_access[i++] = std::log(a); - } - { - T a = 1; - res_access[i++] = std::acos(a); - } - { - T a = 0; - res_access[i++] = std::asin(a); - } - { - T a = 0; - res_access[i++] = std::atan(a); - } - { - T a = 0; - T b = 1; - res_access[i++] = std::atan2(a, b); - } - { - T a = 0; - res_access[i++] = std::cosh(a); - } - { - T a = 0; - res_access[i++] = std::exp(a); - } - { - T a = 1.5; - T b = 1; - res_access[i++] = std::fmod(a, b); - } - { - T a = 0; - res_access[i++] = std::frexp(a, &exp_access[0]); - } - { - T a = 1; - res_access[i++] = std::ldexp(a, 1); - } - { - T a = 1; - res_access[i++] = std::log10(a); - } - { - T a = 1; - res_access[i++] = std::modf(a, &iptr_access[0]); - } - { - T a = 1; - T b = 1; - res_access[i++] = std::pow(a, b); - } - { - T a = 0; - res_access[i++] = std::sinh(a); - } - { - T a = 4; - res_access[i++] = std::sqrt(a); - } - { - T a = 0; - res_access[i++] = std::tan(a); - } - { - T a = 0; - res_access[i++] = std::tanh(a); - } - { - T a = 1; - res_access[i++] = std::acosh(a); - } - { - T a = 0; - res_access[i++] = std::asinh(a); - } - { - T a = 0; - res_access[i++] = std::atanh(a); - } - { - T a = 1; - res_access[i++] = std::cbrt(a); - } - { - T a = 0; - res_access[i++] = std::erf(a); - } - { - T a = 0; - res_access[i++] = std::erfc(a); - } - { - T a = 1; - res_access[i++] = std::exp2(a); - } - { - T a = 0; - res_access[i++] = std::expm1(a); - } - { - T a = 1; - T b = 0; - res_access[i++] = std::fdim(a, b); - } - { - T a = 1; - T b = 1; - T c = 1; - res_access[i++] = std::fma(a, b, c); - } - { - T a = 3; - T b = 4; - res_access[i++] = std::hypot(a, b); - } - { - T a = 1; - res_access[i++] = std::ilogb(a); - } - { - T a = 0; - res_access[i++] = std::log1p(a); - } - { - T a = 1; - res_access[i++] = std::log2(a); - } - { - T a = 1; - res_access[i++] = std::logb(a); - } - { - T a = 0.5; - T b = 1; - res_access[i++] = std::remainder(a, b); - } - { - T a = 0.5; - T b = 1; - res_access[i++] = std::remquo(a, b, &quo_access[0]); - } - { - T a = NAN; - res_access[i++] = std::tgamma(a); - } - { - T a = NAN; - res_access[i++] = std::lgamma(a); - } + res_access[i++] = std::cos(0.0f); + res_access[i++] = std::sin(0.0f); + res_access[i++] = std::log(1.0f); + res_access[i++] = std::acos(1.0f); + res_access[i++] = std::asin(0.0f); + res_access[i++] = std::atan(0.0f); + res_access[i++] = std::atan2(0.0f, 1.0f); + res_access[i++] = std::cosh(0.0f); + res_access[i++] = std::exp(0.0f); + res_access[i++] = std::fmod(1.5f, 1.0f); + res_access[i++] = std::frexp(0.0f, &exp_access[0]); + res_access[i++] = std::ldexp(1.0f, 1); + res_access[i++] = std::log10(1.0f); + res_access[i++] = std::modf(1.0f, &iptr_access[0]); + res_access[i++] = std::pow(1.0f, 1.0f); + res_access[i++] = std::sinh(0.0f); + res_access[i++] = std::sqrt(4.0f); + res_access[i++] = std::tan(0.0f); + res_access[i++] = std::tanh(0.0f); + res_access[i++] = std::acosh(1.0f); + res_access[i++] = std::asinh(0.0f); + res_access[i++] = std::atanh(0.0f); + res_access[i++] = std::cbrt(1.0f); + res_access[i++] = std::erf(0.0f); + res_access[i++] = std::erfc(0.0f); + res_access[i++] = std::exp2(1.0f); + res_access[i++] = std::expm1(0.0f); + res_access[i++] = std::fdim(1.0f, 0.0f); + res_access[i++] = std::fma(1.0f, 1.0f, 1.0f); + res_access[i++] = std::hypot(3.0f, 4.0f); + res_access[i++] = std::ilogb(1.0f); + res_access[i++] = std::log1p(0.0f); + res_access[i++] = std::log2(1.0f); + res_access[i++] = std::logb(1.0f); + res_access[i++] = std::remainder(0.5f, 1.0f); + res_access[i++] = std::remquo(0.5f, 1.0f, &quo_access[0]); + T a = NAN; + res_access[i++] = std::tgamma(a); + res_access[i++] = std::lgamma(a); }); }); } + // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } + // Test modf integral part assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/math_fp64_test.cpp b/sycl/test/devicelib/math_fp64_test.cpp index 5b9eb95695b20..a48cc0905edc3 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -23,10 +23,13 @@ double refIptr = 1; void device_math_test(s::queue &deviceQueue) { s::range<1> numOfItems{TEST_NUM}; double result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function int exponent = -1; + // Variable iptr stores the integral part of float point in modf function double iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { @@ -40,179 +43,61 @@ void device_math_test(s::queue &deviceQueue) { auto iptr_access = buffer3.template get_access(cgh); auto quo_access = buffer4.template get_access(cgh); cgh.single_task([=]() { - int i = 0; - { - double a = 0; - res_access[i++] = cos(a); - } - { - double a = 0; - res_access[i++] = sin(a); - } - { - double a = 1; - res_access[i++] = log(a); - } - { - double a = 1; - res_access[i++] = acos(a); - } - { - double a = 0; - res_access[i++] = asin(a); - } - { - double a = 0; - res_access[i++] = atan(a); - } - { - double a = 0; - double b = 1; - res_access[i++] = atan2(a, b); - } - { - double a = 0; - res_access[i++] = cosh(a); - } - { - double a = 0; - res_access[i++] = exp(a); - } - { - double a = 1.5; - double b = 1; - res_access[i++] = fmod(a, b); - } - { - double a = 0; - res_access[i++] = frexp(a, &exp_access[0]); - } - { - double a = 1; - res_access[i++] = ldexp(a, 1); - } - { - double a = 1; - res_access[i++] = log10(a); - } - { - double a = 1; - res_access[i++] = modf(a, &iptr_access[0]); - } - { - double a = 1; - double b = 1; - res_access[i++] = pow(a, b); - } - { - double a = 0; - res_access[i++] = sinh(a); - } - { - double a = 4; - res_access[i++] = sqrt(a); - } - { - double a = 0; - res_access[i++] = tan(a); - } - { - double a = 0; - res_access[i++] = tanh(a); - } - { - double a = 1; - res_access[i++] = acosh(a); - } - { - double a = 0; - res_access[i++] = asinh(a); - } - { - double a = 0; - res_access[i++] = atanh(a); - } - { - double a = 1; - res_access[i++] = cbrt(a); - } - { - double a = 0; - res_access[i++] = erf(a); - } - { - double a = 0; - res_access[i++] = erfc(a); - } - { - double a = 1; - res_access[i++] = exp2(a); - } - { - double a = 0; - res_access[i++] = expm1(a); - } - { - double a = 1; - double b = 0; - res_access[i++] = fdim(a, b); - } - { - double a = 1; - double b = 1; - double c = 1; - res_access[i++] = fma(a, b, c); - } - { - double a = 3; - double b = 4; - res_access[i++] = hypot(a, b); - } - { - double a = 1; - res_access[i++] = ilogb(a); - } - { - double a = 0; - res_access[i++] = log1p(a); - } - { - double a = 1; - res_access[i++] = log2(a); - } - { - double a = 1; - res_access[i++] = logb(a); - } - { - double a = 0.5; - double b = 1; - res_access[i++] = remainder(a, b); - } - { - double a = 0.5; - double b = 1; - res_access[i++] = remquo(a, b, &quo_access[0]); - } - { - double a = NAN; - res_access[i++] = tgamma(a); - } - { - double a = NAN; - res_access[i++] = lgamma(a); - } + int i = 0; + res_access[i++] = std::cos(0.0); + res_access[i++] = std::sin(0.0); + res_access[i++] = std::log(1.0); + res_access[i++] = std::acos(1.0); + res_access[i++] = std::asin(0.0); + res_access[i++] = std::atan(0.0); + res_access[i++] = std::atan2(0.0, 1.0); + res_access[i++] = std::cosh(0.0); + res_access[i++] = std::exp(0.0); + res_access[i++] = std::fmod(1.5, 1.0); + res_access[i++] = std::frexp(0.0, &exp_access[0]); + res_access[i++] = std::ldexp(1.0, 1); + res_access[i++] = std::log10(1.0); + res_access[i++] = std::modf(1.0, &iptr_access[0]); + res_access[i++] = std::pow(1.0, 1.0); + res_access[i++] = std::sinh(0.0); + res_access[i++] = std::sqrt(4.0); + res_access[i++] = std::tan(0.0); + res_access[i++] = std::tanh(0.0); + res_access[i++] = std::acosh(1.0); + res_access[i++] = std::asinh(0.0); + res_access[i++] = std::atanh(0.0); + res_access[i++] = std::cbrt(1.0); + res_access[i++] = std::erf(0.0); + res_access[i++] = std::erfc(0.0); + res_access[i++] = std::exp2(1.0); + res_access[i++] = std::expm1(0.0); + res_access[i++] = std::fdim(1.0, 0.0); + res_access[i++] = std::fma(1.0, 1.0, 1.0); + res_access[i++] = std::hypot(3.0, 4.0); + res_access[i++] = std::ilogb(1.0); + res_access[i++] = std::log1p(0.0); + res_access[i++] = std::log2(1.0); + res_access[i++] = std::logb(1.0); + res_access[i++] = std::remainder(0.5, 1.0); + res_access[i++] = std::remquo(0.5, 1.0, &quo_access[0]); + double a = NAN; + res_access[i++] = std::tgamma(a); + res_access[i++] = std::lgamma(a); }); }); } + // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } + // Test modf integral part assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index 6ec771c941a09..3550fac1aecd9 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -23,10 +23,13 @@ float refIptr = 1; void device_math_test(s::queue &deviceQueue) { s::range<1> numOfItems{TEST_NUM}; float result[TEST_NUM] = {-1}; + // Variable exponent is an integer value to store the exponent in frexp function int exponent = -1; + // Variable iptr stores the integral part of float point in modf function float iptr = -1; + // Variable quo stores the sign and some bits of x/y in remquo function int quo = -1; { @@ -41,178 +44,60 @@ void device_math_test(s::queue &deviceQueue) { auto quo_access = buffer4.template get_access(cgh); cgh.single_task([=]() { int i = 0; - { - float a = 0; - res_access[i++] = cosf(a); - } - { - float a = 0; - res_access[i++] = sinf(a); - } - { - float a = 1; - res_access[i++] = logf(a); - } - { - float a = 1; - res_access[i++] = acosf(a); - } - { - float a = 0; - res_access[i++] = asinf(a); - } - { - float a = 0; - res_access[i++] = atanf(a); - } - { - float a = 0; - float b = 1; - res_access[i++] = atan2f(a, b); - } - { - float a = 0; - res_access[i++] = coshf(a); - } - { - float a = 0; - res_access[i++] = expf(a); - } - { - float a = 1.5; - float b = 1; - res_access[i++] = fmodf(a, b); - } - { - float a = 0; - res_access[i++] = frexpf(a, &exp_access[0]); - } - { - float a = 1; - res_access[i++] = ldexpf(a, 1); - } - { - float a = 1; - res_access[i++] = log10f(a); - } - { - float a = 1; - res_access[i++] = modff(a, &iptr_access[0]); - } - { - float a = 1; - float b = 1; - res_access[i++] = powf(a, b); - } - { - float a = 0; - res_access[i++] = sinhf(a); - } - { - float a = 4; - res_access[i++] = sqrtf(a); - } - { - float a = 0; - res_access[i++] = tanf(a); - } - { - float a = 0; - res_access[i++] = tanhf(a); - } - { - float a = 1; - res_access[i++] = acoshf(a); - } - { - float a = 0; - res_access[i++] = asinhf(a); - } - { - float a = 0; - res_access[i++] = atanhf(a); - } - { - float a = 1; - res_access[i++] = cbrtf(a); - } - { - float a = 0; - res_access[i++] = erff(a); - } - { - float a = 0; - res_access[i++] = erfcf(a); - } - { - float a = 1; - res_access[i++] = exp2f(a); - } - { - float a = 0; - res_access[i++] = expm1f(a); - } - { - float a = 1; - float b = 0; - res_access[i++] = fdimf(a, b); - } - { - float a = 1; - float b = 1; - float c = 1; - res_access[i++] = fmaf(a, b, c); - } - { - float a = 3; - float b = 4; - res_access[i++] = hypotf(a, b); - } - { - float a = 1; - res_access[i++] = ilogbf(a); - } - { - float a = 0; - res_access[i++] = log1pf(a); - } - { - float a = 1; - res_access[i++] = log2f(a); - } - { - float a = 1; - res_access[i++] = logbf(a); - } - { - float a = 0.5; - float b = 1; - res_access[i++] = remainderf(a, b); - } - { - float a = 0.5; - float b = 1; - res_access[i++] = remquof(a, b, &quo_access[0]); - } - { - float a = NAN; - res_access[i++] = tgammaf(a); - } - { - float a = NAN; - res_access[i++] = lgammaf(a); - } + res_access[i++] = std::cos(0.0f); + res_access[i++] = std::sin(0.0f); + res_access[i++] = std::log(1.0f); + res_access[i++] = std::acos(1.0f); + res_access[i++] = std::asin(0.0f); + res_access[i++] = std::atan(0.0f); + res_access[i++] = std::atan2(0.0f, 1.0f); + res_access[i++] = std::cosh(0.0f); + res_access[i++] = std::exp(0.0f); + res_access[i++] = std::fmod(1.5f, 1.0f); + res_access[i++] = std::frexp(0.0f, &exp_access[0]); + res_access[i++] = std::ldexp(1.0f, 1); + res_access[i++] = std::log10(1.0f); + res_access[i++] = std::modf(1.0f, &iptr_access[0]); + res_access[i++] = std::pow(1.0f, 1.0f); + res_access[i++] = std::sinh(0.0f); + res_access[i++] = std::sqrt(4.0f); + res_access[i++] = std::tan(0.0f); + res_access[i++] = std::tanh(0.0f); + res_access[i++] = std::acosh(1.0f); + res_access[i++] = std::asinh(0.0f); + res_access[i++] = std::atanh(0.0f); + res_access[i++] = std::cbrt(1.0f); + res_access[i++] = std::erf(0.0f); + res_access[i++] = std::erfc(0.0f); + res_access[i++] = std::exp2(1.0f); + res_access[i++] = std::expm1(0.0f); + res_access[i++] = std::fdim(1.0f, 0.0f); + res_access[i++] = std::fma(1.0f, 1.0f, 1.0f); + res_access[i++] = std::hypot(3.0f, 4.0f); + res_access[i++] = std::ilogb(1.0f); + res_access[i++] = std::log1p(0.0f); + res_access[i++] = std::log2(1.0f); + res_access[i++] = std::logb(1.0f); + res_access[i++] = std::remainder(0.5f, 1.0f); + res_access[i++] = std::remquo(0.5f, 1.0f, &quo_access[0]); + float a = NAN; + res_access[i++] = std::tgamma(a); + res_access[i++] = std::lgamma(a); }); }); } + // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { assert(is_about_FP(result[i], ref[i])); } + // Test modf integral part assert(is_about_FP(iptr, refIptr)); + // Test frexp exponent assert(exponent == 0); + // Test remquo sign assert(quo == 0); } diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index d2de1cb508911..e109281357d74 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -4,27 +4,30 @@ // T must be float-point type template bool is_about_FP(T x, T y) { + bool ret; + // At least one input is nan if (std::isnan(x) || std::isnan(y)) - return std::isnan(x) && std::isnan(y); + ret = std::isnan(x) && std::isnan(y); // At least one input is inf else if (std::isinf(x) || std::isinf(y)) - return (x == y); + ret = (x == y); // two finite else { + T threshold = std::numeric_limits::epsilon() * 100; if (x != 0 && y != 0) { T max_v = std::fmax(std::abs(x), std::abs(y)); - return (std::abs(x - y) / max_v) < - std::numeric_limits::epsilon() * 100; + ret = (std::abs(x - y) / max_v) < threshold; } else { if (x != 0) - return std::abs(x) < std::numeric_limits::epsilon() * 100; + ret = std::abs(x) < threshold; else - return std::abs(y) < std::numeric_limits::epsilon() * 100; + ret = std::abs(y) < threshold; } } + return ret; } #endif From fd16ef2016a829f342be9d36e4c5d336d66d7d47 Mon Sep 17 00:00:00 2001 From: haonanya Date: Mon, 16 Mar 2020 17:39:06 +0800 Subject: [PATCH 10/19] [SYCL] [Test] Improve readability for math_utils.hpp Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index e109281357d74..3773ef6585f91 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -8,10 +8,10 @@ bool is_about_FP(T x, T y) { // At least one input is nan if (std::isnan(x) || std::isnan(y)) - ret = std::isnan(x) && std::isnan(y); + return std::isnan(x) && std::isnan(y); // At least one input is inf - else if (std::isinf(x) || std::isinf(y)) + if (std::isinf(x) || std::isinf(y)) ret = (x == y); // two finite From 764518e19826bb77fe500506110f428892e7bd73 Mon Sep 17 00:00:00 2001 From: haonanya <61303033+haonanya@users.noreply.github.com> Date: Wed, 18 Mar 2020 09:30:02 +0800 Subject: [PATCH 11/19] Update sycl/test/devicelib/math_utils.hpp Modify code to follow coding standards Co-Authored-By: Ronan Keryell Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index 3773ef6585f91..7f58d57bf8a8d 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -19,7 +19,7 @@ bool is_about_FP(T x, T y) { T threshold = std::numeric_limits::epsilon() * 100; if (x != 0 && y != 0) { T max_v = std::fmax(std::abs(x), std::abs(y)); - ret = (std::abs(x - y) / max_v) < threshold; + return std::abs(x - y) < threshold*max_v; } else { if (x != 0) ret = std::abs(x) < threshold; From 10c938b57ca0bd3ab4af4f69e148f431f20620f0 Mon Sep 17 00:00:00 2001 From: haonanya Date: Wed, 18 Mar 2020 11:41:27 +0800 Subject: [PATCH 12/19] [SYCL] Add comment for is_about_FP function in math_utils.hpp Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index 7f58d57bf8a8d..e88795f1e1abe 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -1,7 +1,11 @@ #ifndef MATH_UTILS #include #include -// T must be float-point type + +// Since it is not proper to compare float point using operator ==, this +// function measures whether the result of cmath function from kernel is +// close to the reference and machine epsilon is used as threshold in this +// function. T must be float-point type. template bool is_about_FP(T x, T y) { bool ret; @@ -19,7 +23,7 @@ bool is_about_FP(T x, T y) { T threshold = std::numeric_limits::epsilon() * 100; if (x != 0 && y != 0) { T max_v = std::fmax(std::abs(x), std::abs(y)); - return std::abs(x - y) < threshold*max_v; + return std::abs(x - y) < threshold * max_v; } else { if (x != 0) ret = std::abs(x) < threshold; From e8bfe401885de0fce528b8e7f6452f265eec5383 Mon Sep 17 00:00:00 2001 From: haonanya Date: Thu, 19 Mar 2020 11:25:41 +0800 Subject: [PATCH 13/19] [SYCL] [Test] Rename is_about_FP function and remove ret to use early return. Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index e88795f1e1abe..3f4ebf93c0327 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -7,8 +7,7 @@ // close to the reference and machine epsilon is used as threshold in this // function. T must be float-point type. template -bool is_about_FP(T x, T y) { - bool ret; +bool approx_equal_fp(T x, T y) { // At least one input is nan if (std::isnan(x) || std::isnan(y)) @@ -16,7 +15,7 @@ bool is_about_FP(T x, T y) { // At least one input is inf if (std::isinf(x) || std::isinf(y)) - ret = (x == y); + return (x == y); // two finite else { @@ -26,12 +25,11 @@ bool is_about_FP(T x, T y) { return std::abs(x - y) < threshold * max_v; } else { if (x != 0) - ret = std::abs(x) < threshold; + return std::abs(x) < threshold; else - ret = std::abs(y) < threshold; + return std::abs(y) < threshold; } } - return ret; } #endif From 014d3b269742f014d05fff4ad6bee67fa54ddcf6 Mon Sep 17 00:00:00 2001 From: haonanya Date: Thu, 19 Mar 2020 11:28:38 +0800 Subject: [PATCH 14/19] [SYCL] [Test] Rename complex compare function. Signed-off-by: haonanya --- .../devicelib/c99_complex_math_fp64_test.cpp | 20 +++++++------- sycl/test/devicelib/c99_complex_math_test.cpp | 20 +++++++------- .../devicelib/std_complex_math_fp64_test.cpp | 26 +++++++++---------- sycl/test/devicelib/std_complex_math_test.cpp | 26 +++++++++---------- 4 files changed, 46 insertions(+), 46 deletions(-) diff --git a/sycl/test/devicelib/c99_complex_math_fp64_test.cpp b/sycl/test/devicelib/c99_complex_math_fp64_test.cpp index b7da3b6a9e37a..c039025b11115 100644 --- a/sycl/test/devicelib/c99_complex_math_fp64_test.cpp +++ b/sycl/test/devicelib/c99_complex_math_fp64_test.cpp @@ -9,8 +9,8 @@ #define CMPLX(r, i) ((double __complex__){ (double)r, (double)i }) #endif -bool is_about_C99_CMPLX(double __complex__ x, double __complex__ y) { - return is_about_FP(creal(x), creal(y)) && is_about_FP(cimag(x), cimag(y)); +bool approx_equal_c99_cmplx(double __complex__ x, double __complex__ y) { + return approx_equal_fp(creal(x), creal(y)) && approx_equal_fp(cimag(x), cimag(y)); } namespace s = cl::sycl; @@ -44,7 +44,7 @@ void device_c99_complex_times(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } @@ -81,7 +81,7 @@ void device_c99_complex_divides(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 8; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } @@ -107,7 +107,7 @@ void device_c99_complex_sqrt(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } @@ -132,7 +132,7 @@ void device_c99_complex_abs(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_FP(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_fp(buf_out2[idx], ref_results2[idx])); } } @@ -158,7 +158,7 @@ void device_c99_complex_exp(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } @@ -184,7 +184,7 @@ void device_c99_complex_log(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } @@ -208,7 +208,7 @@ void device_c99_complex_sin(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } @@ -232,7 +232,7 @@ void device_c99_complex_cos(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_C99_CMPLX(buf_out2[idx], ref_results2[idx])); + assert(approx_equal_c99_cmplx(buf_out2[idx], ref_results2[idx])); } } diff --git a/sycl/test/devicelib/c99_complex_math_test.cpp b/sycl/test/devicelib/c99_complex_math_test.cpp index 9637ccd4a0568..704d80bd0131f 100644 --- a/sycl/test/devicelib/c99_complex_math_test.cpp +++ b/sycl/test/devicelib/c99_complex_math_test.cpp @@ -10,8 +10,8 @@ #define CMPLXF(r, i) ((float __complex__){ (float)r, (float)i }) #endif -bool is_about_C99_CMPLXF(float __complex__ x, float __complex__ y) { - return is_about_FP(crealf(x), crealf(y)) && is_about_FP(cimagf(x), cimagf(y)); +bool approx_equal_c99_cmplxf(float __complex__ x, float __complex__ y) { + return approx_equal_fp(crealf(x), crealf(y)) && approx_equal_fp(cimagf(x), cimagf(y)); } namespace s = cl::sycl; @@ -46,7 +46,7 @@ void device_c99_complex_times(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } @@ -83,7 +83,7 @@ void device_c99_complex_divides(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 8; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } @@ -110,7 +110,7 @@ void device_c99_complex_sqrt(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } @@ -136,7 +136,7 @@ void device_c99_complex_abs(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_FP(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_fp(buf_out1[idx], ref_results1[idx])); } } @@ -162,7 +162,7 @@ void device_c99_complex_exp(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } @@ -188,7 +188,7 @@ void device_c99_complex_log(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } @@ -212,7 +212,7 @@ void device_c99_complex_sin(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } @@ -236,7 +236,7 @@ void device_c99_complex_cos(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_C99_CMPLXF(buf_out1[idx], ref_results1[idx])); + assert(approx_equal_c99_cmplxf(buf_out1[idx], ref_results1[idx])); } } diff --git a/sycl/test/devicelib/std_complex_math_fp64_test.cpp b/sycl/test/devicelib/std_complex_math_fp64_test.cpp index 31e0f65a388fc..74a83604342be 100644 --- a/sycl/test/devicelib/std_complex_math_fp64_test.cpp +++ b/sycl/test/devicelib/std_complex_math_fp64_test.cpp @@ -11,8 +11,8 @@ constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; template -bool is_about_CMPLX(complex x, complex y) { - return is_about_FP(x.real(), y.real()) && is_about_FP(x.imag(), y.imag()); +bool approx_equal_cmplx(complex x, complex y) { + return approx_equal_fp(x.real(), y.real()) && approx_equal_fp(x.imag(), y.imag()); } template @@ -45,7 +45,7 @@ void device_complex_times(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -84,7 +84,7 @@ void device_complex_divides(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 8; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -112,7 +112,7 @@ void device_complex_sqrt(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -139,7 +139,7 @@ void device_complex_norm(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_FP(buf_out[idx], ref_results[idx])); + assert(approx_equal_fp(buf_out[idx], ref_results[idx])); } } @@ -167,7 +167,7 @@ void device_complex_abs(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_FP(buf_out[idx], ref_results[idx])); + assert(approx_equal_fp(buf_out[idx], ref_results[idx])); } } @@ -195,7 +195,7 @@ void device_complex_exp(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -223,7 +223,7 @@ void device_complex_log(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -248,7 +248,7 @@ void device_complex_log10(s::queue &deviceQueue) { }); } - assert(is_about_CMPLX(buf_out, ref_result)); + assert(approx_equal_cmplx(buf_out, ref_result)); } template @@ -273,7 +273,7 @@ void device_complex_sin(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -299,7 +299,7 @@ void device_complex_cos(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -326,7 +326,7 @@ void device_complex_polar(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } diff --git a/sycl/test/devicelib/std_complex_math_test.cpp b/sycl/test/devicelib/std_complex_math_test.cpp index af4ef799fa2f4..f71cfcb7f7088 100644 --- a/sycl/test/devicelib/std_complex_math_test.cpp +++ b/sycl/test/devicelib/std_complex_math_test.cpp @@ -12,8 +12,8 @@ constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; template -bool is_about_CMPLX(complex x, complex y) { - return is_about_FP(x.real(), y.real()) && is_about_FP(x.imag(), y.imag()); +bool approx_equal_cmplx(complex x, complex y) { + return approx_equal_fp(x.real(), y.real()) && approx_equal_fp(x.imag(), y.imag()); } template @@ -46,7 +46,7 @@ void device_complex_times(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -85,7 +85,7 @@ void device_complex_divides(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 8; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -113,7 +113,7 @@ void device_complex_sqrt(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -140,7 +140,7 @@ void device_complex_norm(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_FP(buf_out[idx], ref_results[idx])); + assert(approx_equal_fp(buf_out[idx], ref_results[idx])); } } @@ -168,7 +168,7 @@ void device_complex_abs(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_FP(buf_out[idx], ref_results[idx])); + assert(approx_equal_fp(buf_out[idx], ref_results[idx])); } } @@ -196,7 +196,7 @@ void device_complex_exp(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -224,7 +224,7 @@ void device_complex_log(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -249,7 +249,7 @@ void device_complex_log10(s::queue &deviceQueue) { }); } - assert(is_about_CMPLX(buf_out, ref_result)); + assert(approx_equal_cmplx(buf_out, ref_result)); } template @@ -274,7 +274,7 @@ void device_complex_sin(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -300,7 +300,7 @@ void device_complex_cos(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 2; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } @@ -327,7 +327,7 @@ void device_complex_polar(s::queue &deviceQueue) { } for (size_t idx = 0; idx < 4; ++idx) { - assert(is_about_CMPLX(buf_out[idx], ref_results[idx])); + assert(approx_equal_cmplx(buf_out[idx], ref_results[idx])); } } From 6ef6faca1e37d951085a1a6a6625d545388fe7ec Mon Sep 17 00:00:00 2001 From: haonanya Date: Thu, 19 Mar 2020 11:32:18 +0800 Subject: [PATCH 15/19] [SYCL] [Test] There is a rename for float point compare in math_utils.hpp, Modify corresponding call in these files. Signed-off-by: haonanya --- sycl/test/devicelib/cmath_fp64_test.cpp | 4 ++-- sycl/test/devicelib/cmath_test.cpp | 4 ++-- sycl/test/devicelib/math_fp64_test.cpp | 4 ++-- sycl/test/devicelib/math_test.cpp | 4 ++-- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/test/devicelib/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp index 294660ca8fc46..8ac2a634623e8 100644 --- a/sycl/test/devicelib/cmath_fp64_test.cpp +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -90,11 +90,11 @@ void device_cmath_test(s::queue &deviceQueue) { // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { - assert(is_about_FP(result[i], ref[i])); + assert(approx_equal_fp(result[i], ref[i])); } // Test modf integral part - assert(is_about_FP(iptr, refIptr)); + assert(approx_equal_fp(iptr, refIptr)); // Test frexp exponent assert(exponent == 0); diff --git a/sycl/test/devicelib/cmath_test.cpp b/sycl/test/devicelib/cmath_test.cpp index 4624b4bccf346..33861bacd3c7a 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -90,11 +90,11 @@ void device_cmath_test(s::queue &deviceQueue) { // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { - assert(is_about_FP(result[i], ref[i])); + assert(approx_equal_fp(result[i], ref[i])); } // Test modf integral part - assert(is_about_FP(iptr, refIptr)); + assert(approx_equal_fp(iptr, refIptr)); // Test frexp exponent assert(exponent == 0); diff --git a/sycl/test/devicelib/math_fp64_test.cpp b/sycl/test/devicelib/math_fp64_test.cpp index a48cc0905edc3..f262932b5fe5a 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -89,11 +89,11 @@ void device_math_test(s::queue &deviceQueue) { // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { - assert(is_about_FP(result[i], ref[i])); + assert(approx_equal_fp(result[i], ref[i])); } // Test modf integral part - assert(is_about_FP(iptr, refIptr)); + assert(approx_equal_fp(iptr, refIptr)); // Test frexp exponent assert(exponent == 0); diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index 3550fac1aecd9..6212273092575 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -89,11 +89,11 @@ void device_math_test(s::queue &deviceQueue) { // Compare result with reference for (int i = 0; i < TEST_NUM; ++i) { - assert(is_about_FP(result[i], ref[i])); + assert(approx_equal_fp(result[i], ref[i])); } // Test modf integral part - assert(is_about_FP(iptr, refIptr)); + assert(approx_equal_fp(iptr, refIptr)); // Test frexp exponent assert(exponent == 0); From 88fc2546fd9e4011fe52b1692fef5a38573cc0db Mon Sep 17 00:00:00 2001 From: haonanya Date: Tue, 7 Apr 2020 21:18:44 +0800 Subject: [PATCH 16/19] Fix use else after return issue Signed-off-by: haonanya --- sycl/test/devicelib/math_utils.hpp | 16 +++++----------- 1 file changed, 5 insertions(+), 11 deletions(-) diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index 3f4ebf93c0327..e48476a0c225e 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -18,18 +18,12 @@ bool approx_equal_fp(T x, T y) { return (x == y); // two finite - else { - T threshold = std::numeric_limits::epsilon() * 100; - if (x != 0 && y != 0) { - T max_v = std::fmax(std::abs(x), std::abs(y)); - return std::abs(x - y) < threshold * max_v; - } else { - if (x != 0) - return std::abs(x) < threshold; - else - return std::abs(y) < threshold; - } + T threshold = std::numeric_limits::epsilon() * 100; + if (x != 0 && y != 0) { + T max_v = std::fmax(std::abs(x), std::abs(y)); + return std::abs(x - y) < threshold * max_v; } + return x != 0 ? std::abs(x) < threshold : std::abs(y) < threshold; } #endif From 048e5d09fd6908cc88af3b9cf92e94ad07b35e8c Mon Sep 17 00:00:00 2001 From: haonanya Date: Thu, 9 Apr 2020 14:30:34 +0800 Subject: [PATCH 17/19] [SYCL][Test] Enable cmath lit test. Signed-off-by: haonanya --- sycl/test/devicelib/cmath_fp64_test.cpp | 3 +++ sycl/test/devicelib/cmath_test.cpp | 3 +++ sycl/test/devicelib/math_fp64_test.cpp | 3 +++ sycl/test/devicelib/math_override_test.cpp | 3 +++ sycl/test/devicelib/math_test.cpp | 3 +++ 5 files changed, 15 insertions(+) diff --git a/sycl/test/devicelib/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp index 8ac2a634623e8..30954e0eff59f 100644 --- a/sycl/test/devicelib/cmath_fp64_test.cpp +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -1,6 +1,9 @@ // UNSUPPORTED: windows // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath-fp64.o -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include #include diff --git a/sycl/test/devicelib/cmath_test.cpp b/sycl/test/devicelib/cmath_test.cpp index 33861bacd3c7a..bb2e37345d6fc 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -1,6 +1,9 @@ // UNSUPPORTED: windows // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath.o -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include #include diff --git a/sycl/test/devicelib/math_fp64_test.cpp b/sycl/test/devicelib/math_fp64_test.cpp index f262932b5fe5a..45fef78494bcd 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -1,6 +1,9 @@ // UNSUPPORTED: windows // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath-fp64.o -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include #include diff --git a/sycl/test/devicelib/math_override_test.cpp b/sycl/test/devicelib/math_override_test.cpp index 62166b95ab1ab..3420734710685 100644 --- a/sycl/test/devicelib/math_override_test.cpp +++ b/sycl/test/devicelib/math_override_test.cpp @@ -1,6 +1,9 @@ // UNSUPPORTED: windows // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath.o -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include #include diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index 6212273092575..0d77a2251caba 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -1,6 +1,9 @@ // UNSUPPORTED: windows // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath.o -o %t.out +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include #include From e8ac45edb91c2230db1174a078fd3a6e97d3efe8 Mon Sep 17 00:00:00 2001 From: haonanya Date: Thu, 9 Apr 2020 20:07:49 +0800 Subject: [PATCH 18/19] [SYCL][Test] Enable Host device for math_override.cpp. Signed-off-by: haonanya --- sycl/test/devicelib/math_override_test.cpp | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/sycl/test/devicelib/math_override_test.cpp b/sycl/test/devicelib/math_override_test.cpp index 3420734710685..bd5068a18f084 100644 --- a/sycl/test/devicelib/math_override_test.cpp +++ b/sycl/test/devicelib/math_override_test.cpp @@ -2,11 +2,10 @@ // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %sycl_libs_dir/libsycl-cmath.o -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include #include +#include "math_utils.hpp" namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; constexpr s::access::mode sycl_write = s::access::mode::write; @@ -14,9 +13,7 @@ constexpr s::access::mode sycl_write = s::access::mode::write; // Dummy function provided by user to override device library // version. SYCL_EXTERNAL -extern "C" float sinf(float x) { - return x + 100; -} +extern "C" float sinf(float x) { return x + 100.f; } class DeviceTest; @@ -34,13 +31,13 @@ void device_test() { cgh.single_task([=]() { // Should use the sin function defined by user, device // library version should be ignored here - res_access_sin[0] = sinf(0); - res_access_cos[0] = cosf(0); + res_access_sin[0] = sinf(0.f); + res_access_cos[0] = cosf(0.f); }); }); } - assert(((int)result_sin == 100) && ((int)result_cos == 1)); + assert(approx_equal_fp(result_sin, 100.f) && approx_equal_fp(result_cos, 1.f)); } int main() { @@ -48,3 +45,4 @@ int main() { std::cout << "Pass" << std::endl; return 0; } + From a54c5f1ea7389e987611df92cbc8a6a10e979b78 Mon Sep 17 00:00:00 2001 From: haonanya Date: Thu, 9 Apr 2020 20:20:25 +0800 Subject: [PATCH 19/19] [SYCL] [Test] Fix format issue. Signed-off-by: haonanya --- sycl/test/devicelib/math_override_test.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/devicelib/math_override_test.cpp b/sycl/test/devicelib/math_override_test.cpp index bd5068a18f084..e634bf77fbd37 100644 --- a/sycl/test/devicelib/math_override_test.cpp +++ b/sycl/test/devicelib/math_override_test.cpp @@ -5,6 +5,7 @@ #include #include #include + #include "math_utils.hpp" namespace s = cl::sycl; constexpr s::access::mode sycl_read = s::access::mode::read; @@ -45,4 +46,3 @@ int main() { std::cout << "Pass" << std::endl; return 0; } -