[SYCL][Test] Add test for libdevice float/double to bfloat16 utils (#9000)

Previously, we added utils function in libdevice to convert double/float
to bfloat16, this PR aims to add corresponding test.

---------

Signed-off-by: jinge90 <[email protected]>

Author: jinge90

sycl/test-e2e/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<sycl::info::device::name>()
+            << "\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<double> 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<uint16_t> 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;
+}

sycl/test-e2e/DeviceLib/imf_float2bfloat16.cpp

@@ -0,0 +1,81 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// RUN: %clangxx -fsycl -fno-builtin -fsycl-device-lib-jit-link %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+//
+// UNSUPPORTED: cuda || hip
+
+// All __imf_* bf16 functions are implemented via fp32 emulation, so we don't
+// need to check whether underlying device supports bf16 or not.
+#include "imf_utils.hpp"
+
+extern "C" {
+uint16_t __imf_float2bfloat16(float);
+uint16_t __imf_float2bfloat16_rd(float);
+uint16_t __imf_float2bfloat16_rn(float);
+uint16_t __imf_float2bfloat16_ru(float);
+uint16_t __imf_float2bfloat16_rz(float);
+};
+
+int main() {
+  sycl::queue device_queue(sycl::default_selector_v);
+  std::cout << "Running on "
+            << device_queue.get_device().get_info<sycl::info::device::name>()
+            << "\n";
+
+  {
+    std::initializer_list<float> input_vals = {
+        __builtin_bit_cast(float, 0x0),        // +0
+        __builtin_bit_cast(float, 0x80000000), // -0
+        __builtin_bit_cast(float, 0x1),        // min positive subnormal
+        __builtin_bit_cast(float, 0x7FFFFF),   // max positive subnormal
+        __builtin_bit_cast(float, 0x5A6BFC),   // positive subnormal
+        __builtin_bit_cast(float, 0x80000001), // max negative subnormal
+        __builtin_bit_cast(float, 0x807FFFFF), // min negative subnormal
+        __builtin_bit_cast(float, 0x805A6FED), // negative subnormal
+        __builtin_bit_cast(float, 0x7F800000), // +inf
+        __builtin_bit_cast(float, 0xFF800000), // -inf
+        __builtin_bit_cast(float, 0x2E05CBA9), // positive normal
+        __builtin_bit_cast(float, 0x7E5A8935), // positive normal
+        __builtin_bit_cast(float, 0xAE4411FC), // negative normal
+        __builtin_bit_cast(float, 0xFA84C773), // negative normal
+        __builtin_bit_cast(float, 0x7F7FFFFF), // max positive normal
+        __builtin_bit_cast(float, 0x765FCEED), // positive normal
+        __builtin_bit_cast(float, 0xFF7FFFFF), // min negative normal
+        __builtin_bit_cast(float, 0xAC763561), // negative normal
+    };
+
+    std::initializer_list<uint16_t> ref_vals = {
+        0x0,    0x8000, 0x0,    0x80,   0x5a,   0x8000, 0x8080, 0x805A, 0x7F80,
+        0xFF80, 0x2E06, 0x7E5B, 0xAE44, 0xFA85, 0x7F80, 0x7660, 0xFF80, 0xAC76};
+
+    std::initializer_list<uint16_t> ref_vals_rd = {
+        0x0,    0x8000, 0x0,    0x7F,   0x5A,   0x8001, 0x8080, 0x805B, 0x7F80,
+        0xFF80, 0x2E05, 0x7E5A, 0xAE45, 0xFA85, 0x7F7F, 0x765F, 0xFF80, 0xAC77};
+
+    std::initializer_list<uint16_t> ref_vals_ru = {
+        0x0,    0x8000, 0x1,    0x80,   0x5B,   0x8000, 0x807F, 0x805A, 0x7F80,
+        0xFF80, 0x2E06, 0x7E5B, 0xAE44, 0xFA84, 0x7F80, 0x7660, 0xFF7F, 0xAC76};
+
+    std::initializer_list<uint16_t> ref_vals_rz = {
+        0x0,    0x8000, 0x0,    0x7F,   0x5A,   0x8000, 0x807F, 0x805A, 0x7F80,
+        0xFF80, 0x2E05, 0x7E5A, 0xAE44, 0xFA84, 0x7F7F, 0x765F, 0xFF7F, 0xAC76};
+
+    test_host(input_vals, ref_vals, F(__imf_float2bfloat16));
+    test_host(input_vals, ref_vals_rd, F(__imf_float2bfloat16_rd));
+    test_host(input_vals, ref_vals, F(__imf_float2bfloat16_rn));
+    test_host(input_vals, ref_vals_ru, F(__imf_float2bfloat16_ru));
+    test_host(input_vals, ref_vals_rz, F(__imf_float2bfloat16_rz));
+    test(device_queue, input_vals, ref_vals, F(__imf_float2bfloat16));
+    test(device_queue, input_vals, ref_vals_rd, F(__imf_float2bfloat16_rd));
+    test(device_queue, input_vals, ref_vals, F(__imf_float2bfloat16_rn));
+    test(device_queue, input_vals, ref_vals_ru, F(__imf_float2bfloat16_ru));
+    test(device_queue, input_vals, ref_vals_rz, F(__imf_float2bfloat16_rz));
+  }
+  return 0;
+}