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/cmath_fp64_test.cpp b/sycl/test/devicelib/cmath_fp64_test.cpp new file mode 100644 index 0000000000000..30954e0eff59f --- /dev/null +++ b/sycl/test/devicelib/cmath_fp64_test.cpp @@ -0,0 +1,116 @@ +// 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 +#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; + +#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{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(&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) { + auto res_access = buffer1.template get_access(cgh); + 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; + 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(approx_equal_fp(result[i], ref[i])); + } + + // Test modf integral part + assert(approx_equal_fp(iptr, refIptr)); + + // Test frexp exponent + assert(exponent == 0); + + // Test remquo sign + assert(quo == 0); +} + +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..bb2e37345d6fc 100644 --- a/sycl/test/devicelib/cmath_test.cpp +++ b/sycl/test/devicelib/cmath_test.cpp @@ -1,79 +1,109 @@ // 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 +#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; +#define TEST_NUM 38 -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); - }); - }); - } +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,}; - assert(result == 1); -} +float refIptr = 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); - }); - }); - } +void device_cmath_test(s::queue &deviceQueue) { + s::range<1> numOfItems{TEST_NUM}; + T result[TEST_NUM] = {-1}; - assert(result == 0); -} + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; -template -class DeviceLog; + // Variable iptr stores the integral part of float point in modf function + T iptr = -1; -template -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - T result = -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 buffer1(result, numOfItems); + 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) { auto res_access = buffer1.template get_access(cgh); - cgh.single_task >([=]() { - res_access[0] = std::log(1); + 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; + 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); }); }); } - assert(result == 0); -} + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { + assert(approx_equal_fp(result[i], ref[i])); + } -template -void device_cmath_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); + // Test modf integral part + assert(approx_equal_fp(iptr, refIptr)); + + // Test frexp exponent + assert(exponent == 0); + + // Test remquo sign + assert(quo == 0); } 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..45fef78494bcd 100644 --- a/sycl/test/devicelib/math_fp64_test.cpp +++ b/sycl/test/devicelib/math_fp64_test.cpp @@ -1,72 +1,108 @@ // 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 +#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 DeviceSin; +#define TEST_NUM 38 -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result_d = -1; - { - s::buffer buffer1(&result_d, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = sin(0); - }); - }); - } +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,}; - assert(result_d == 0); -} - -class DeviceCos; +double refIptr = 1; -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result_d = -1; - { - s::buffer buffer1(&result_d, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = cos(0); - }); - }); - } +void device_math_test(s::queue &deviceQueue) { + s::range<1> numOfItems{TEST_NUM}; + double result[TEST_NUM] = {-1}; - assert(result_d == 1); -} + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; -class DeviceLog; + // Variable iptr stores the integral part of float point in modf function + double iptr = -1; -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - double result_d = -1; + // Variable quo stores the sign and some bits of x/y in remquo function + int quo = -1; { - s::buffer buffer2(&result_d, numOfItems); + s::buffer buffer1(result, numOfItems); + 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) { - auto res_access1 = buffer2.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = log(1); + auto res_access = buffer1.template get_access(cgh); + 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; + 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); }); }); } - assert(result_d == 0); -} + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { + assert(approx_equal_fp(result[i], ref[i])); + } -void device_math_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); + // Test modf integral part + assert(approx_equal_fp(iptr, refIptr)); + + // Test frexp exponent + assert(exponent == 0); + + // Test remquo sign + assert(quo == 0); } int main() { diff --git a/sycl/test/devicelib/math_override_test.cpp b/sycl/test/devicelib/math_override_test.cpp index 62166b95ab1ab..e634bf77fbd37 100644 --- a/sycl/test/devicelib/math_override_test.cpp +++ b/sycl/test/devicelib/math_override_test.cpp @@ -1,9 +1,12 @@ // 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 #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; @@ -11,9 +14,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; @@ -31,13 +32,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() { diff --git a/sycl/test/devicelib/math_test.cpp b/sycl/test/devicelib/math_test.cpp index 1b5aa4332fe8e..0d77a2251caba 100644 --- a/sycl/test/devicelib/math_test.cpp +++ b/sycl/test/devicelib/math_test.cpp @@ -1,72 +1,108 @@ // 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 +#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 DeviceSin; +#define TEST_NUM 38 -void device_sin_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result_f = -1; - { - s::buffer buffer1(&result_f, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = sinf(0); - }); - }); - } +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,}; - assert(result_f == 0); -} - -class DeviceCos; +float refIptr = 1; -void device_cos_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result_f = -1; - { - s::buffer buffer1(&result_f, numOfItems); - deviceQueue.submit([&](cl::sycl::handler &cgh) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = cosf(0); - }); - }); - } +void device_math_test(s::queue &deviceQueue) { + s::range<1> numOfItems{TEST_NUM}; + float result[TEST_NUM] = {-1}; - assert(result_f == 1); -} + // Variable exponent is an integer value to store the exponent in frexp function + int exponent = -1; -class DeviceLog; + // Variable iptr stores the integral part of float point in modf function + float iptr = -1; -void device_log_test(s::queue &deviceQueue) { - s::range<1> numOfItems{1}; - float result_f = -1; + // Variable quo stores the sign and some bits of x/y in remquo function + int quo = -1; { - s::buffer buffer1(&result_f, numOfItems); + s::buffer buffer1(result, numOfItems); + 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) { - auto res_access1 = buffer1.get_access(cgh); - cgh.single_task([=]() { - res_access1[0] = logf(1); + auto res_access = buffer1.template get_access(cgh); + 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; + 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); }); }); } - assert(result_f == 0); -} + // Compare result with reference + for (int i = 0; i < TEST_NUM; ++i) { + assert(approx_equal_fp(result[i], ref[i])); + } -void device_math_test(s::queue &deviceQueue) { - device_cos_test(deviceQueue); - device_sin_test(deviceQueue); - device_log_test(deviceQueue); + // Test modf integral part + assert(approx_equal_fp(iptr, refIptr)); + + // Test frexp exponent + assert(exponent == 0); + + // Test remquo sign + assert(quo == 0); } int main() { diff --git a/sycl/test/devicelib/math_utils.hpp b/sycl/test/devicelib/math_utils.hpp index ef13d36a54ba9..e48476a0c225e 100644 --- a/sycl/test/devicelib/math_utils.hpp +++ b/sycl/test/devicelib/math_utils.hpp @@ -1,25 +1,29 @@ #ifndef MATH_UTILS #include #include -using namespace std; -// 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) { - if (x == 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; - } - else { - if (x != 0) - return abs(x) < numeric_limits::epsilon() * 100; - else - return abs(y) < numeric_limits::epsilon() * 100; - } +bool approx_equal_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 + if (std::isinf(x) || std::isinf(y)) + return (x == y); + + // two finite + 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 diff --git a/sycl/test/devicelib/std_complex_math_fp64_test.cpp b/sycl/test/devicelib/std_complex_math_fp64_test.cpp index e1213a713ad5b..74a83604342be 100644 --- a/sycl/test/devicelib/std_complex_math_fp64_test.cpp +++ b/sycl/test/devicelib/std_complex_math_fp64_test.cpp @@ -4,14 +4,15 @@ #include #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; 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 @@ -44,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])); } } @@ -83,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])); } } @@ -111,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])); } } @@ -138,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])); } } @@ -166,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])); } } @@ -194,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])); } } @@ -222,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])); } } @@ -247,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 @@ -272,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])); } } @@ -298,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])); } } @@ -325,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 f3fe5ae9ae510..f71cfcb7f7088 100644 --- a/sycl/test/devicelib/std_complex_math_test.cpp +++ b/sycl/test/devicelib/std_complex_math_test.cpp @@ -5,13 +5,15 @@ #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; 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 @@ -44,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])); } } @@ -83,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])); } } @@ -111,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])); } } @@ -138,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])); } } @@ -166,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])); } } @@ -194,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])); } } @@ -222,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])); } } @@ -247,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 @@ -272,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])); } } @@ -298,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])); } } @@ -325,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])); } }