diff --git a/sycl/test-e2e/DeviceLib/built-ins/offload-prec-div-sqrt.cpp b/sycl/test-e2e/DeviceLib/built-ins/offload-prec-div-sqrt.cpp new file mode 100644 index 0000000000000..e00e123ca1aaf --- /dev/null +++ b/sycl/test-e2e/DeviceLib/built-ins/offload-prec-div-sqrt.cpp @@ -0,0 +1,56 @@ +// RUN: %{build} -foffload-fp32-prec-div -foffload-fp32-prec-sqrt -o %t.out +// RUN: %{run} %t.out + +// Test if div and sqrt become precise from IEEE-754 perspective when +// -foffload-fp32-prec-div -foffload-fp32-prec-sqrt are passed. + +#include +#include +#include + +constexpr float value = 560.0f; +constexpr float divider = 279.9f; + +void test_div() { + sycl::queue q(sycl::default_selector_v); + float *inValue = (float *)sycl::malloc_shared(sizeof(float), q); + float *inDivider = (float *)sycl::malloc_shared(sizeof(float), q); + float *output = (float *)sycl::malloc_shared(sizeof(float), q); + *inValue = value; + *inDivider = divider; + q.submit([&](sycl::handler &h) { + h.single_task([=] { + float res = *inValue / *inDivider; + *output = res; + }); + }).wait(); + + float hostRef = value / divider; + int ulpDist = std::abs(sycl::bit_cast(hostRef) - + sycl::bit_cast(*output)); + assert(ulpDist == 0 && "Division is not precise"); +} + +void test_sqrt() { + sycl::queue q(sycl::default_selector_v); + float *inValue = (float *)sycl::malloc_shared(sizeof(float), q); + float *output = (float *)sycl::malloc_shared(sizeof(float), q); + *inValue = value; + q.submit([&](sycl::handler &h) { + h.single_task([=] { + float res = sycl::sqrt(*inValue); + *output = res; + }); + }).wait(); + + float hostRef = std::sqrt(value); + int ulpDist = std::abs(sycl::bit_cast(hostRef) - + sycl::bit_cast(*output)); + assert(ulpDist == 0 && "Sqrt is not precise"); +} + +int main() { + test_div(); + test_sqrt(); + return 0; +}