Add max & min bfloat16 support & tests

Also reimplement `fmax_nan` & `fmin_nan` in terms of improved max/min

Author: joeatodd

sycl/include/syclcompat/math.hpp

@@ -732,7 +732,7 @@ cbrt(ValueT val) {
 // For floating-point types, `float` or `double` arguments are acceptable.
 // For integer types, `std::uint32_t`, `std::int32_t`, `std::uint64_t` or
 // `std::int64_t` type arguments are acceptable.
-// sycl::half supported as well.
+// sycl::half supported as well, and sycl::ext::oneapi::bfloat16 if available.
 template <typename ValueT, typename ValueU>
 inline std::enable_if_t<std::is_integral_v<ValueT> &&
                             std::is_integral_v<ValueU>,
@@ -741,15 +741,23 @@ min(ValueT a, ValueU b) {
   return sycl::min(static_cast<std::common_type_t<ValueT, ValueU>>(a),
                    static_cast<std::common_type_t<ValueT, ValueU>>(b));
 }
+
 template <typename ValueT, typename ValueU>
-inline std::enable_if_t<std::is_floating_point_v<ValueT> &&
-                            std::is_floating_point_v<ValueU>,
+inline std::enable_if_t<syclcompat::is_floating_point_v<ValueT> &&
+                            syclcompat::is_floating_point_v<ValueU>,
                         std::common_type_t<ValueT, ValueU>>
 min(ValueT a, ValueU b) {
-  return sycl::fmin(static_cast<std::common_type_t<ValueT, ValueU>>(a),
-                    static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  if constexpr (detail::support_bfloat16_math &&
+                std::is_same_v<std::common_type_t<ValueT, ValueU>,
+                               sycl::ext::oneapi::bfloat16>) {
+    return sycl::ext::oneapi::experimental::fmin(
+        static_cast<std::common_type_t<ValueT, ValueU>>(a),
+        static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  } else {
+    return sycl::fmin(static_cast<std::common_type_t<ValueT, ValueU>>(a),
+                      static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  }
 }
-inline sycl::half min(sycl::half a, sycl::half b) { return sycl::fmin(a, b); }
 
 template <typename ValueT, typename ValueU>
 inline std::enable_if_t<std::is_integral_v<ValueT> &&
@@ -760,14 +768,21 @@ max(ValueT a, ValueU b) {
                    static_cast<std::common_type_t<ValueT, ValueU>>(b));
 }
 template <typename ValueT, typename ValueU>
-inline std::enable_if_t<std::is_floating_point_v<ValueT> &&
-                            std::is_floating_point_v<ValueU>,
+inline std::enable_if_t<syclcompat::is_floating_point_v<ValueT> &&
+                            syclcompat::is_floating_point_v<ValueU>,
                         std::common_type_t<ValueT, ValueU>>
 max(ValueT a, ValueU b) {
-  return sycl::fmax(static_cast<std::common_type_t<ValueT, ValueU>>(a),
-                    static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  if constexpr (detail::support_bfloat16_math &&
+                std::is_same_v<std::common_type_t<ValueT, ValueU>,
+                               sycl::ext::oneapi::bfloat16>) {
+    return sycl::ext::oneapi::experimental::fmax(
+        static_cast<std::common_type_t<ValueT, ValueU>>(a),
+        static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  } else {
+    return sycl::fmax(static_cast<std::common_type_t<ValueT, ValueU>>(a),
+                      static_cast<std::common_type_t<ValueT, ValueU>>(b));
+  }
 }
-inline sycl::half max(sycl::half a, sycl::half b) { return sycl::fmax(a, b); }
 
 /// Performs 2 elements comparison and returns the bigger one. If either of
 /// inputs is NaN, then return NaN.
@@ -779,16 +794,7 @@ inline std::common_type_t<ValueT, ValueU> fmax_nan(const ValueT a,
                                                    const ValueU b) {
   if (detail::isnan(a) || detail::isnan(b))
     return NAN;
-  if constexpr (detail::support_bfloat16_math &&
-                std::is_same_v<std::common_type_t<ValueT, ValueU>,
-                               sycl::ext::oneapi::bfloat16>) {
-    return sycl::ext::oneapi::experimental::fmax(
-        static_cast<std::common_type_t<ValueT, ValueU>>(a),
-        static_cast<std::common_type_t<ValueT, ValueU>>(b));
-  } else {
-    return sycl::fmax(static_cast<std::common_type_t<ValueT, ValueU>>(a),
-                      static_cast<std::common_type_t<ValueT, ValueU>>(b));
-  }
+  return syclcompat::max(a, b);
 }
 
 template <typename ValueT, typename ValueU>
@@ -813,16 +819,7 @@ inline std::common_type_t<ValueT, ValueU> fmin_nan(const ValueT a,
                                                    const ValueU b) {
   if (detail::isnan(a) || detail::isnan(b))
     return NAN;
-  if constexpr (detail::support_bfloat16_math &&
-                std::is_same_v<std::common_type_t<ValueT, ValueU>,
-                               sycl::ext::oneapi::bfloat16>) {
-    return sycl::ext::oneapi::experimental::fmin(
-        static_cast<std::common_type_t<ValueT, ValueU>>(a),
-        static_cast<std::common_type_t<ValueT, ValueU>>(b));
-  } else {
-    return sycl::fmin(static_cast<std::common_type_t<ValueT, ValueU>>(a),
-                      static_cast<std::common_type_t<ValueT, ValueU>>(b));
-  }
+  return syclcompat::min(a,b);
 }
 
 template <typename ValueT, typename ValueU>

sycl/test-e2e/syclcompat/math/math_ops.cpp

@@ -34,7 +34,7 @@
 template <typename ValueT, typename ValueU>
 inline void max_kernel(ValueT *a, ValueU *b,
                        std::common_type_t<ValueT, ValueU> *r) {
-  *r = syclcompat::max(*a, *b);
+  *r = syclcompat::max<ValueT, ValueU>(*a, *b);
 }
 
 template <typename ValueT, typename ValueU = ValueT>
@@ -54,7 +54,7 @@ void test_syclcompat_max() {
 template <typename ValueT, typename ValueU>
 inline void min_kernel(ValueT *a, ValueU *b,
                        std::common_type_t<ValueT, ValueU> *r) {
-  *r = syclcompat::min(*a, *b);
+  *r = syclcompat::min<ValueT,ValueU>(*a, *b);
 }
 
 template <typename ValueT, typename ValueU = ValueT>
@@ -342,8 +342,15 @@ int main() {
   // Basic testing of deduction to avoid combinatorial explosion
   test_syclcompat_max<double, float>();
   test_syclcompat_max<long, int>();
+#ifdef SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS
+  test_syclcompat_max<sycl::ext::oneapi::bfloat16, float>();
+#endif
+
   test_syclcompat_min<double, float>();
   test_syclcompat_min<long, int>();
+#ifdef SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS
+  test_syclcompat_min<sycl::ext::oneapi::bfloat16, float>();
+#endif
 
   INSTANTIATE_ALL_TYPES(fp_type_list, test_syclcompat_fmin_nan);
   INSTANTIATE_ALL_CONTAINER_TYPES(fp_type_list, sycl::vec, test_container_syclcompat_fmin_nan);