diff --git a/libdevice/imf_rounding_op.hpp b/libdevice/imf_rounding_op.hpp index f76c32990ac44..4bbf15ef94855 100644 --- a/libdevice/imf_rounding_op.hpp +++ b/libdevice/imf_rounding_op.hpp @@ -670,12 +670,24 @@ template Ty __fp_div(Ty x, Ty y, int rd) { const UTy sig_off_mask = (one_bits << (sizeof(UTy) * 8 - 1)) - 1; if (((x_exp == __iml_fp_config::exp_mask) && (x_fra != 0x0)) || - ((y_exp == __iml_fp_config::exp_mask) && (y_fra != 0x0)) || - ((y_bit & sig_off_mask) == 0x0)) { + ((y_exp == __iml_fp_config::exp_mask) && (y_fra != 0x0))) { UTy tmp = __iml_fp_config::nan_bits; return __builtin_bit_cast(Ty, tmp); } + // 0.f / 0.f ----> NAN + if ((y_bit & sig_off_mask) == 0x0) { + if ((x_bit & sig_off_mask) == 0x0) { + UTy tmp = __iml_fp_config::nan_bits; + return __builtin_bit_cast(Ty, tmp); + } else { + // return +inf if x_sig and y_sig are same otherwise return -inf + UTy tmp = (z_sig == 0) ? __iml_fp_config::pos_inf_bits + : __iml_fp_config::neg_inf_bits; + return __builtin_bit_cast(Ty, tmp); + } + } + if ((x_exp == __iml_fp_config::exp_mask) && (x_fra == 0x0)) { if ((y_exp == __iml_fp_config::exp_mask) && (y_fra == 0x0)) { UTy tmp = __iml_fp_config::nan_bits; diff --git a/sycl/test-e2e/DeviceLib/imf/a.out b/sycl/test-e2e/DeviceLib/imf/a.out new file mode 100755 index 0000000000000..87c1920b1b2a5 Binary files /dev/null and b/sycl/test-e2e/DeviceLib/imf/a.out differ diff --git a/sycl/test-e2e/DeviceLib/imf/fp32_rounding_test.cpp b/sycl/test-e2e/DeviceLib/imf/fp32_rounding_test.cpp index e11b05c9aa6c4..55621e72920bf 100644 --- a/sycl/test-e2e/DeviceLib/imf/fp32_rounding_test.cpp +++ b/sycl/test-e2e/DeviceLib/imf/fp32_rounding_test.cpp @@ -92,18 +92,24 @@ int main(int, char **) { } { - std::initializer_list input_vals1 = {0x1p-1, 0x1.8bd054p+6, - 0x1.fcd686p+0, -0x1.7f9abp+3}; - std::initializer_list input_vals2 = {-0x1.a8p+2, -0x1.674a3cp+5, - 0x1.f3d6aep+10, 0x1.d6bf48p+10}; - std::initializer_list ref_vals_rd = {0xbd9a90e8, 0xc00d030d, - 0x3a824df9, 0xbbd09c3a}; - std::initializer_list ref_vals_rn = {0xbd9a90e8, 0xc00d030c, - 0x3a824df9, 0xbbd09c39}; - std::initializer_list ref_vals_ru = {0xbd9a90e7, 0xc00d030c, - 0x3a824dfa, 0xbbd09c39}; - std::initializer_list ref_vals_rz = {0xbd9a90e7, 0xc00d030c, - 0x3a824df9, 0xbbd09c39}; + std::initializer_list input_vals1 = { + 0x1p-1, 0x1.8bd054p+6, 0x1.fcd686p+0, -0x1.7f9abp+3, + 0x1p+0, -0x1p+0, 0x0p+0}; + std::initializer_list input_vals2 = { + -0x1.a8p+2, -0x1.674a3cp+5, 0x1.f3d6aep+10, 0x1.d6bf48p+10, + 0x0p+0, 0x0p+0, 0x0p+0}; + std::initializer_list ref_vals_rd = { + 0xbd9a90e8, 0xc00d030d, 0x3a824df9, 0xbbd09c3a, + 0x7F800000, 0xFF800000, 0x7FC00000}; + std::initializer_list ref_vals_rn = { + 0xbd9a90e8, 0xc00d030c, 0x3a824df9, 0xbbd09c39, + 0x7F800000, 0xFF800000, 0x7FC00000}; + std::initializer_list ref_vals_ru = { + 0xbd9a90e7, 0xc00d030c, 0x3a824dfa, 0xbbd09c39, + 0x7F800000, 0xFF800000, 0x7FC00000}; + std::initializer_list ref_vals_rz = { + 0xbd9a90e7, 0xc00d030c, 0x3a824df9, 0xbbd09c39, + 0x7F800000, 0xFF800000, 0x7FC00000}; test2(device_queue, input_vals1, input_vals2, ref_vals_rd, F2T(uint32_t, sycl::ext::intel::math::fdiv_rd)); std::cout << "sycl::ext::intel::math::fdiv_rd passes." << std::endl; diff --git a/sycl/test-e2e/DeviceLib/imf/fp64_rounding_test.cpp b/sycl/test-e2e/DeviceLib/imf/fp64_rounding_test.cpp index a07d8e3c4d4bd..bdcbf49534188 100644 --- a/sycl/test-e2e/DeviceLib/imf/fp64_rounding_test.cpp +++ b/sycl/test-e2e/DeviceLib/imf/fp64_rounding_test.cpp @@ -115,24 +115,36 @@ int main(int, char **) { } { - std::initializer_list input_vals1 = { - 0x1.5ef3da7bf609ap+4, 0x1.fbd37afb0f8edp-1, 0x1.9238e38e38e35p+6, - 0x1.7p+3}; - std::initializer_list input_vals2 = { - -0x1.bc7db6de6d33fp+9, 0x1.2f638fa4e71a6p+10, 0x1.08e38e38e38e3p+4, - -0x1.94p+3}; + std::initializer_list input_vals1 = {0x1.5ef3da7bf609ap+4, + 0x1.fbd37afb0f8edp-1, + 0x1.9238e38e38e35p+6, + 0x1.7p+3, + 0x1p+0, + -0x1p+0, + 0x0p+0}; + std::initializer_list input_vals2 = {-0x1.bc7db6de6d33fp+9, + 0x1.2f638fa4e71a6p+10, + 0x1.08e38e38e38e3p+4, + -0x1.94p+3, + 0x0p+0, + 0x0p+0, + 0x0p+0}; std::initializer_list ref_vals_rd = { 0xbf994414312c26ab, 0x3f4ac811fc63acd9, 0x40184b98e9aa180a, - 0xbfed260511be1959}; + 0xbfed260511be1959, 0x7FF0000000000000, 0xFFF0000000000000, + 0x7FF8000000000000}; std::initializer_list ref_vals_rn = { 0xbf994414312c26ab, 0x3f4ac811fc63acd9, 0x40184b98e9aa180b, - 0xbfed260511be1959}; + 0xbfed260511be1959, 0x7FF0000000000000, 0xFFF0000000000000, + 0x7FF8000000000000}; std::initializer_list ref_vals_ru = { 0xbf994414312c26aa, 0x3f4ac811fc63acda, 0x40184b98e9aa180b, - 0xbfed260511be1958}; + 0xbfed260511be1958, 0x7FF0000000000000, 0xFFF0000000000000, + 0x7FF8000000000000}; std::initializer_list ref_vals_rz = { 0xbf994414312c26aa, 0x3f4ac811fc63acd9, 0x40184b98e9aa180a, - 0xbfed260511be1958}; + 0xbfed260511be1958, 0x7FF0000000000000, 0xFFF0000000000000, + 0x7FF8000000000000}; test2(device_queue, input_vals1, input_vals2, ref_vals_rd, F2T(uint64_t, sycl::ext::intel::math::ddiv_rd)); std::cout << "sycl::ext::intel::math::ddiv_rd passes." << std::endl; diff --git a/sycl/test-e2e/DeviceLib/imf/imf_utils.hpp b/sycl/test-e2e/DeviceLib/imf/imf_utils.hpp index b6f43f30514a0..4fe9f4bfdaed2 100644 --- a/sycl/test-e2e/DeviceLib/imf/imf_utils.hpp +++ b/sycl/test-e2e/DeviceLib/imf/imf_utils.hpp @@ -14,6 +14,11 @@ typedef _Float16 _iml_half_internal; typedef uint16_t _iml_half_internal; #endif +#define EXP_MASK32 0xFF +#define EXP_MASK64 0x7FF +#define FRA_MASK32 0x7FFFFF +#define FRA_MASK64 0xFFFFFFFFFFFFF + template class imf_utils_default_equ { public: bool operator()(Ty x, Ty y) { @@ -24,6 +29,28 @@ template class imf_utils_default_equ { }; }; +template <> class imf_utils_default_equ { +public: + bool operator()(uint32_t x, uint32_t y) { + bool x_is_nan = + (((x >> 23) & EXP_MASK32) == EXP_MASK32) && ((x & FRA_MASK32) != 0); + bool y_is_nan = + (((y >> 23) & EXP_MASK32) == EXP_MASK32) && ((y & FRA_MASK32) != 0); + return (x_is_nan && y_is_nan) || (x == y); + } +}; + +template <> class imf_utils_default_equ { +public: + bool operator()(uint64_t x, uint64_t y) { + bool x_is_nan = + (((x >> 52) & EXP_MASK64) == EXP_MASK64) && ((x & FRA_MASK64) != 0); + bool y_is_nan = + (((y >> 52) & EXP_MASK64) == EXP_MASK64) && ((y & FRA_MASK64) != 0); + return (x_is_nan && y_is_nan) || (x == y); + } +}; + // Used to test half precision utils template >