diff --git a/SYCL/DeviceLib/imf_double2bfloat16.cpp b/SYCL/DeviceLib/imf_double2bfloat16.cpp new file mode 100644 index 0000000000..4971409795 --- /dev/null +++ b/SYCL/DeviceLib/imf_double2bfloat16.cpp @@ -0,0 +1,64 @@ +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t1.out +// RUN: %GPU_RUN_PLACEHOLDER %t1.out +// +// UNSUPPORTED: cuda || hip + +#include "imf_utils.hpp" + +extern "C" { +uint16_t __imf_double2bfloat16(double); +} + +int main() { + + sycl::queue device_queue(sycl::default_selector_v); + std::cout << "Running on " + << device_queue.get_device().get_info() + << "\n"; + + if (!device_queue.get_device().has(sycl::aspect::fp64)) { + std::cout << "Test skipped on platform without fp64 support." << std::endl; + return 0; + } + + { + std::initializer_list input_vals = { + __builtin_bit_cast(double, 0ULL), // 0 + __builtin_bit_cast(double, 0x7FF0000000000000), // +infinity + __builtin_bit_cast(double, 0xFFF0000000000000), // -infinity + __builtin_bit_cast(double, 0x4026800000000000), // 11.25 + __builtin_bit_cast(double, 0x409025643C8E4F03), // 1033.3478872524 + __builtin_bit_cast(double, 0x40EFFC0000000000), // 65504 + __builtin_bit_cast(double, 0xC0EFFC0000000000), // -65504 + __builtin_bit_cast(double, 0xC0D38814311F5D54), // -20000.31549820055245 + __builtin_bit_cast(double, 0x409F9B8D12ACEFA7), // 2022.887766554 + __builtin_bit_cast(double, 0x40ee120000000000), // 61584 + __builtin_bit_cast(double, 0xC0EE160000000000), // -61616 + __builtin_bit_cast(double, 0x40FAA93000000000), // 109203 + __builtin_bit_cast(double, 0xC1A7D8B7FF20E365), // -200039423.564234872 + __builtin_bit_cast(double, 0x3C370EF54646D497), // 1.25e-18 + __builtin_bit_cast(double, 0xBCB1B3CFC61ACF52), // -2.4567e-16 + __builtin_bit_cast(double, 0x39F036448D68D482), // 1.2789e-29 + __builtin_bit_cast(double, 0xB99C100A89BE0A2D), // -3.45899e-31 + __builtin_bit_cast(double, 0x47EFFFFFFFFFFFFF), + __builtin_bit_cast(double, 0x47EFF00000000000), + __builtin_bit_cast(double, 0x47EFD00000000000), + __builtin_bit_cast(double, 0xC7EFFFFFFFFFFFFF), + __builtin_bit_cast(double, 0xC7EFF00000000000), + __builtin_bit_cast(double, 0xC7EFD00000000000), + __builtin_bit_cast(double, 0x37AFFFFFFFFFFFFF), + __builtin_bit_cast(double, 0x380FFFFFFFFFFFFF), + }; + std::initializer_list ref_vals = { + 0x0, 0x7F80, 0xFF80, 0x4134, 0x4481, 0x4780, 0xC780, 0xC69C, 0x44FD, + 0x4771, 0xC771, 0x47D5, 0xCD3F, 0x21B8, 0xA58E, 0x0F82, 0x8CE1, 0x7F80, + 0x7F80, 0x7F7E, 0xFF80, 0xFF80, 0xFF7E, 0x2, 0x80}; + test_host(input_vals, ref_vals, F(__imf_double2bfloat16)); + test(device_queue, input_vals, ref_vals, F(__imf_double2bfloat16)); + } + + return 0; +}