[SYCL][Test] Enable bf16 math test on all backends (#8999)

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

Author: jinge90

sycl/include/sycl/ext/oneapi/bfloat16.hpp

@@ -9,12 +9,9 @@
 #pragma once
 
 #include <CL/__spirv/spirv_ops.hpp>
+#include <sycl/builtins.hpp>
 #include <sycl/half_type.hpp>
 
-#if !defined(__SYCL_DEVICE_ONLY__)
-#include <cmath>
-#endif
-
 extern "C" __DPCPP_SYCL_EXTERNAL uint16_t
 __devicelib_ConvertFToBF16INTEL(const float &) noexcept;
 extern "C" __DPCPP_SYCL_EXTERNAL float
@@ -46,15 +43,8 @@ class bfloat16 {
   ~bfloat16() = default;
 
 private:
-  // Explicit conversion functions
-  static detail::Bfloat16StorageT from_float(const float &a) {
-#if defined(__SYCL_DEVICE_ONLY__)
-#if defined(__NVPTX__)
-#if (__SYCL_CUDA_ARCH__ >= 800)
-    return __nvvm_f2bf16_rn(a);
-#else
-    // TODO find a better way to check for NaN
-    if (a != a)
+  static detail::Bfloat16StorageT from_float_fallback(const float &a) {
+    if (sycl::isnan(a))
       return 0xffc1;
     union {
       uint32_t intStorage;
@@ -64,23 +54,24 @@ class bfloat16 {
     // Do RNE and truncate
     uint32_t roundingBias = ((intStorage >> 16) & 0x1) + 0x00007FFF;
     return static_cast<uint16_t>((intStorage + roundingBias) >> 16);
+  }
+
+  // Explicit conversion functions
+  static detail::Bfloat16StorageT from_float(const float &a) {
+#if defined(__SYCL_DEVICE_ONLY__)
+#if defined(__NVPTX__)
+#if (__SYCL_CUDA_ARCH__ >= 800)
+    return __nvvm_f2bf16_rn(a);
+#else
+    return from_float_fallback(a);
 #endif
+#elif defined(__AMDGCN__)
+    return from_float_fallback(a);
 #else
     return __devicelib_ConvertFToBF16INTEL(a);
 #endif
-#else
-    // In case float value is nan - propagate bfloat16's qnan
-    if (std::isnan(a))
-      return 0xffc1;
-    union {
-      uint32_t intStorage;
-      float floatValue;
-    };
-    floatValue = a;
-    // Do RNE and truncate
-    uint32_t roundingBias = ((intStorage >> 16) & 0x1) + 0x00007FFF;
-    return static_cast<uint16_t>((intStorage + roundingBias) >> 16);
 #endif
+    return from_float_fallback(a);
   }
 
   static float to_float(const detail::Bfloat16StorageT &a) {

sycl/test-e2e/BFloat16/bfloat16_builtins.cpp

@@ -1,11 +1,7 @@
-// REQUIRES: cuda
-//
-// Currently this test fails to compile for backends other than cuda.
-// Other backends could use this test when bfloat16 math function support is
-// added.
-//
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -Xsycl-target-backend --cuda-gpu-arch=sm_80
-// RUN: %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %if cuda %{ -Xsycl-target-backend --cuda-gpu-arch=sm_80 %} %s -o %t.out
+// Currently the feature isn't supported on FPGA.
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
 #include <sycl/sycl.hpp>
 
 #include <cmath>
@@ -40,7 +36,8 @@ bool check(bool a, bool b) { return (a != b); }
                                                                      cgh);     \
       accessor<int, 1, access::mode::write, target::device> ERR(err_buf, cgh); \
       cgh.parallel_for(N, [=](id<1> index) {                                   \
-        if (check(NAME(bfloat16{A[index]}), NAME(A[index]))) {                 \
+        if (check(sycl::ext::oneapi::experimental::NAME(bfloat16{A[index]}),   \
+                  sycl::NAME(A[index]))) {                                     \
           ERR[0] = 1;                                                          \
         }                                                                      \
       });                                                                      \
@@ -63,7 +60,7 @@ bool check(bool a, bool b) { return (a != b); }
         }                                                                      \
         marray<RETTY, SZ> res = NAME(arg);                                     \
         for (int i = 0; i < SZ; i++) {                                         \
-          if (check(res[i], NAME(A[index][i]))) {                              \
+          if (check(res[i], sycl::NAME(A[index][i]))) {                        \
             ERR[0] = 1;                                                        \
           }                                                                    \
         }                                                                      \