From 1e61fe317ef21534f7b0fc37e13582a181fe85ec Mon Sep 17 00:00:00 2001
From: Chris Perkins <chris.perkins@intel.com>
Date: Mon, 11 Dec 2023 19:43:54 -0800
Subject: [PATCH] Bfloat16 support for sycl::vec

---
 .../sycl/detail/generic_type_lists.hpp        |  62 ++++++++--
 .../sycl/detail/generic_type_traits.hpp       |  17 ++-
 sycl/include/sycl/ext/oneapi/bfloat16.hpp     |  63 +++++++++-
 .../oneapi/matrix/matrix-unified-utils.hpp    |   3 +
 sycl/include/sycl/stream.hpp                  |  39 +++++-
 sycl/include/sycl/types.hpp                   | 111 ++++++++++++++----
 sycl/test/basic_tests/generic_type_traits.cpp |   4 +
 .../implicit_device_copyable_types.cpp        |   3 +
 sycl/test/basic_tests/types.cpp               |  12 ++
 9 files changed, 264 insertions(+), 50 deletions(-)

diff --git a/sycl/include/sycl/detail/generic_type_lists.hpp b/sycl/include/sycl/detail/generic_type_lists.hpp
index a18bb5fd764f..476f33e1cb3a 100644
--- a/sycl/include/sycl/detail/generic_type_lists.hpp
+++ b/sycl/include/sycl/detail/generic_type_lists.hpp
@@ -12,6 +12,8 @@
 #include <sycl/detail/type_list.hpp> // for type_list, address_space_list
 #include <sycl/half_type.hpp>        // for half
 
+#include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16
+
 #include <cstddef>     // for byte, size_t
 #include <type_traits> // for conditional_t, is_signed_v, is_...
 
@@ -41,6 +43,28 @@ using scalar_vector_half_list = tl_append<scalar_half_list, vector_half_list>;
 using half_list =
     tl_append<scalar_half_list, vector_half_list, marray_half_list>;
 
+using scalar_bfloat16_list = type_list<sycl::ext::oneapi::bfloat16>;
+
+using vector_bfloat16_list = type_list<
+    vec<sycl::ext::oneapi::bfloat16, 1>, vec<sycl::ext::oneapi::bfloat16, 2>,
+    vec<sycl::ext::oneapi::bfloat16, 3>, vec<sycl::ext::oneapi::bfloat16, 4>,
+    vec<sycl::ext::oneapi::bfloat16, 8>, vec<sycl::ext::oneapi::bfloat16, 16>>;
+
+using marray_bfloat16_list = type_list<marray<sycl::ext::oneapi::bfloat16, 1>,
+                                       marray<sycl::ext::oneapi::bfloat16, 2>,
+                                       marray<sycl::ext::oneapi::bfloat16, 3>,
+                                       marray<sycl::ext::oneapi::bfloat16, 4>,
+                                       marray<sycl::ext::oneapi::bfloat16, 8>,
+                                       marray<sycl::ext::oneapi::bfloat16, 16>>;
+
+using scalar_vector_bfloat16_list =
+    tl_append<scalar_bfloat16_list, vector_bfloat16_list>;
+
+using bfloat16_list =
+    tl_append<scalar_bfloat16_list, vector_bfloat16_list, marray_bfloat16_list>;
+
+using half_bfloat16_list = tl_append<scalar_half_list, scalar_bfloat16_list>;
+
 using scalar_float_list = type_list<float>;
 
 using vector_float_list =
@@ -73,14 +97,14 @@ using scalar_vector_double_list =
 using double_list =
     tl_append<scalar_double_list, vector_double_list, marray_double_list>;
 
-using scalar_floating_list =
-    tl_append<scalar_float_list, scalar_double_list, scalar_half_list>;
+using scalar_floating_list = tl_append<scalar_float_list, scalar_double_list,
+                                       scalar_half_list, scalar_bfloat16_list>;
 
-using vector_floating_list =
-    tl_append<vector_float_list, vector_double_list, vector_half_list>;
+using vector_floating_list = tl_append<vector_float_list, vector_double_list,
+                                       vector_half_list, vector_bfloat16_list>;
 
-using marray_floating_list =
-    tl_append<marray_float_list, marray_double_list, marray_half_list>;
+using marray_floating_list = tl_append<marray_float_list, marray_double_list,
+                                       marray_half_list, marray_bfloat16_list>;
 
 using scalar_vector_floating_list =
     tl_append<scalar_floating_list, vector_floating_list>;
@@ -91,6 +115,8 @@ using floating_list =
 // geometric floating point types
 using scalar_geo_half_list = type_list<half>;
 
+using scalar_geo_bfloat16_list = type_list<sycl::ext::oneapi::bfloat16>;
+
 using scalar_geo_float_list = type_list<float>;
 
 using scalar_geo_double_list = type_list<double>;
@@ -98,6 +124,10 @@ using scalar_geo_double_list = type_list<double>;
 using vector_geo_half_list =
     type_list<vec<half, 1>, vec<half, 2>, vec<half, 3>, vec<half, 4>>;
 
+using vector_geo_bfloat16_list = type_list<
+    vec<sycl::ext::oneapi::bfloat16, 1>, vec<sycl::ext::oneapi::bfloat16, 2>,
+    vec<sycl::ext::oneapi::bfloat16, 3>, vec<sycl::ext::oneapi::bfloat16, 4>>;
+
 using vector_geo_float_list =
     type_list<vec<float, 1>, vec<float, 2>, vec<float, 3>, vec<float, 4>>;
 
@@ -112,16 +142,21 @@ using marray_geo_double_list =
 
 using geo_half_list = tl_append<scalar_geo_half_list, vector_geo_half_list>;
 
+using geo_bfloat16_list =
+    tl_append<scalar_geo_bfloat16_list, vector_geo_bfloat16_list>;
+
 using geo_float_list = tl_append<scalar_geo_float_list, vector_geo_float_list>;
 
 using geo_double_list =
     tl_append<scalar_geo_double_list, vector_geo_double_list>;
 
-using scalar_geo_list = tl_append<scalar_geo_half_list, scalar_geo_float_list,
-                                  scalar_geo_double_list>;
+using scalar_geo_list =
+    tl_append<scalar_geo_half_list, scalar_geo_bfloat16_list,
+              scalar_geo_float_list, scalar_geo_double_list>;
 
-using vector_geo_list = tl_append<vector_geo_half_list, vector_geo_float_list,
-                                  vector_geo_double_list>;
+using vector_geo_list =
+    tl_append<vector_geo_half_list, vector_geo_bfloat16_list,
+              vector_geo_float_list, vector_geo_double_list>;
 
 using marray_geo_list =
     tl_append<marray_geo_float_list, marray_geo_double_list>;
@@ -131,12 +166,15 @@ using geo_list = tl_append<scalar_geo_list, vector_geo_list>;
 // cross floating point types
 using cross_half_list = type_list<vec<half, 3>, vec<half, 4>>;
 
+using cross_bfloat16_list = type_list<vec<sycl::ext::oneapi::bfloat16, 3>,
+                                      vec<sycl::ext::oneapi::bfloat16, 4>>;
+
 using cross_float_list = type_list<vec<float, 3>, vec<float, 4>>;
 
 using cross_double_list = type_list<vec<double, 3>, vec<double, 4>>;
 
-using cross_floating_list =
-    tl_append<cross_float_list, cross_double_list, cross_half_list>;
+using cross_floating_list = tl_append<cross_float_list, cross_double_list,
+                                      cross_half_list, cross_bfloat16_list>;
 
 using cross_marray_list = type_list<marray<float, 3>, marray<float, 4>,
                                     marray<double, 3>, marray<double, 4>>;
diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp
index 7e8b6a272df8..fea829057d35 100644
--- a/sycl/include/sycl/detail/generic_type_traits.hpp
+++ b/sycl/include/sycl/detail/generic_type_traits.hpp
@@ -56,6 +56,14 @@ inline constexpr bool is_genfloath_v = is_contained_v<T, gtl::half_list>;
 template <typename T>
 inline constexpr bool is_half_v = is_contained_v<T, gtl::scalar_half_list>;
 
+template <typename T>
+inline constexpr bool is_bfloat16_v =
+    is_contained_v<T, gtl::scalar_bfloat16_list>;
+
+template <typename T>
+inline constexpr bool is_half_or_bf16_v =
+    is_contained_v<T, gtl::half_bfloat16_list>;
+
 template <typename T>
 inline constexpr bool is_svgenfloath_v =
     is_contained_v<T, gtl::scalar_vector_half_list>;
@@ -539,10 +547,9 @@ using select_cl_scalar_t = std::conditional_t<
     std::is_integral_v<T>, select_cl_scalar_integral_t<T>,
     std::conditional_t<
         std::is_floating_point_v<T>, select_cl_scalar_float_t<T>,
-        // half is a special case: it is implemented differently on
-        // host and device and therefore, might lower to different
-        // types
-        std::conditional_t<is_half_v<T>,
+        // half and bfloat16 are special cases: they are implemented differently
+        // on host and device and therefore might lower to different types
+        std::conditional_t<is_half_or_bf16_v<T>,
                            sycl::detail::half_impl::BIsRepresentationT,
                            select_cl_scalar_complex_or_T_t<T>>>>;
 
@@ -559,7 +566,7 @@ struct select_cl_vector_or_scalar_or_ptr<
       // select_cl_scalar_t returns _Float16, so, we try to instantiate vec
       // class with _Float16 DataType, which is not expected there
       // So, leave vector<half, N> as-is
-      vec<std::conditional_t<is_half_v<mptr_or_vec_elem_type_t<T>>,
+      vec<std::conditional_t<is_half_or_bf16_v<mptr_or_vec_elem_type_t<T>>,
                              mptr_or_vec_elem_type_t<T>,
                              select_cl_scalar_t<mptr_or_vec_elem_type_t<T>>>,
           T::size()>;
diff --git a/sycl/include/sycl/ext/oneapi/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/bfloat16.hpp
index bd3052e9a048..9fdc0e7c9fac 100644
--- a/sycl/include/sycl/ext/oneapi/bfloat16.hpp
+++ b/sycl/include/sycl/ext/oneapi/bfloat16.hpp
@@ -9,7 +9,6 @@
 #pragma once
 
 #include <sycl/aliases.hpp>                   // for half
-#include <sycl/builtins.hpp>                  // for isnan
 #include <sycl/detail/defines_elementary.hpp> // for __DPCPP_SYCL_EXTERNAL
 #include <sycl/half_type.hpp>                 // for half
 
@@ -22,6 +21,13 @@ __devicelib_ConvertBF16ToFINTEL(const uint16_t &) noexcept;
 
 namespace sycl {
 inline namespace _V1 {
+
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+// forward declaration of sycl::isnan built-in.
+// extern __DPCPP_SYCL_EXTERNAL bool isnan(float a);
+bool isnan(float a);
+#endif
+
 namespace ext::oneapi {
 
 class bfloat16;
@@ -30,9 +36,31 @@ namespace detail {
 using Bfloat16StorageT = uint16_t;
 Bfloat16StorageT bfloat16ToBits(const bfloat16 &Value);
 bfloat16 bitsToBfloat16(const Bfloat16StorageT Value);
+
+// sycl::vec support
+namespace bf16 {
+#ifdef __SYCL_DEVICE_ONLY__
+using Vec2StorageT = Bfloat16StorageT __attribute__((ext_vector_type(2)));
+using Vec3StorageT = Bfloat16StorageT __attribute__((ext_vector_type(3)));
+using Vec4StorageT = Bfloat16StorageT __attribute__((ext_vector_type(4)));
+using Vec8StorageT = Bfloat16StorageT __attribute__((ext_vector_type(8)));
+using Vec16StorageT = Bfloat16StorageT __attribute__((ext_vector_type(16)));
+#else
+using Vec2StorageT = std::array<Bfloat16StorageT, 2>;
+using Vec3StorageT = std::array<Bfloat16StorageT, 3>;
+using Vec4StorageT = std::array<Bfloat16StorageT, 4>;
+using Vec8StorageT = std::array<Bfloat16StorageT, 8>;
+using Vec16StorageT = std::array<Bfloat16StorageT, 16>;
+#endif
+} // namespace bf16
+
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+static bool float_is_nan(float x) { return x != x; }
+#endif
 } // namespace detail
 
 class bfloat16 {
+protected:
   detail::Bfloat16StorageT value;
 
   friend inline detail::Bfloat16StorageT
@@ -42,13 +70,21 @@ class bfloat16 {
 
 public:
   bfloat16() = default;
-  bfloat16(const bfloat16 &) = default;
+  constexpr bfloat16(const bfloat16 &) = default;
+  constexpr bfloat16(bfloat16 &&) = default;
+  constexpr bfloat16 &operator=(const bfloat16 &rhs) = default;
   ~bfloat16() = default;
 
 private:
   static detail::Bfloat16StorageT from_float_fallback(const float &a) {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
     if (sycl::isnan(a))
       return 0xffc1;
+#else
+    if (detail::float_is_nan(a))
+      return 0xffc1;
+#endif
+
     union {
       uint32_t intStorage;
       float floatValue;
@@ -92,6 +128,14 @@ class bfloat16 {
 #endif
   }
 
+protected:
+  friend class sycl::vec<bfloat16, 1>;
+  friend class sycl::vec<bfloat16, 2>;
+  friend class sycl::vec<bfloat16, 3>;
+  friend class sycl::vec<bfloat16, 4>;
+  friend class sycl::vec<bfloat16, 8>;
+  friend class sycl::vec<bfloat16, 16>;
+
 public:
   // Implicit conversion from float to bfloat16
   bfloat16(const float &a) { value = from_float(a); }
@@ -128,7 +172,7 @@ class bfloat16 {
 #elif defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__)
     return bfloat16{-__devicelib_ConvertBF16ToFINTEL(lhs.value)};
 #else
-    return -to_float(lhs.value);
+    return bfloat16{-to_float(lhs.value)};
 #endif
   }
 
@@ -199,6 +243,19 @@ class bfloat16 {
 
   // Bitwise(|,&,~,^), modulo(%) and shift(<<,>>) operations are not supported
   // for floating-point types.
+
+  // Stream Operator << and >>
+  inline friend std::ostream &operator<<(std::ostream &O, bfloat16 const &rhs) {
+    O << static_cast<float>(rhs);
+    return O;
+  }
+
+  inline friend std::istream &operator>>(std::istream &I, bfloat16 &rhs) {
+    float ValFloat = 0.0f;
+    I >> ValFloat;
+    rhs = ValFloat;
+    return I;
+  }
 };
 
 namespace detail {
diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp
index 8a50c435fb0c..4695d461805b 100644
--- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp
+++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified-utils.hpp
@@ -7,6 +7,9 @@
 // ===--------------------------------------------------------------------=== //
 
 #pragma once
+
+#include <optional>
+
 namespace sycl {
 inline namespace _V1 {
 namespace ext {
diff --git a/sycl/include/sycl/stream.hpp b/sycl/include/sycl/stream.hpp
index dd79bd532cc6..915f67ef063d 100644
--- a/sycl/include/sycl/stream.hpp
+++ b/sycl/include/sycl/stream.hpp
@@ -20,6 +20,7 @@
 #include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
 #include <sycl/detail/item_base.hpp>          // for id, range
 #include <sycl/detail/owner_less_base.hpp>    // for OwnerLessBase
+#include <sycl/ext/oneapi/bfloat16.hpp>       // for bfloat16
 #include <sycl/group.hpp>                     // for group
 #include <sycl/h_item.hpp>                    // for h_item
 #include <sycl/half_type.hpp>                 // for half, operator-, operator<
@@ -83,10 +84,10 @@ constexpr size_t MAX_ARRAY_SIZE =
 constexpr unsigned FLUSH_BUF_OFFSET_SIZE = 2;
 
 template <class F, class T = void>
-using EnableIfFP = typename std::enable_if_t<std::is_same_v<F, float> ||
-                                                 std::is_same_v<F, double> ||
-                                                 std::is_same_v<F, half>,
-                                             T>;
+using EnableIfFP = typename std::enable_if_t<
+    std::is_same_v<F, float> || std::is_same_v<F, double> ||
+        std::is_same_v<F, half> || std::is_same_v<F, ext::oneapi::bfloat16>,
+    T>;
 
 using GlobalBufAccessorT = accessor<char, 1, sycl::access::mode::read_write,
                                     sycl::access::target::device>;
@@ -346,6 +347,26 @@ checkForInfNan(char *Buf, T Val) {
   return 0;
 }
 
+template <typename T>
+inline typename std::enable_if_t<std::is_same_v<T, ext::oneapi::bfloat16>,
+                                 unsigned>
+checkForInfNan(char *Buf, T Val) {
+  if (Val != Val)
+    return append(Buf, "nan");
+
+  // Extract the sign from the bits
+  const uint16_t Sign = reinterpret_cast<uint16_t &>(Val) & 0x8000;
+  // Extract the exponent from the bits
+  const uint16_t Exp16 = (reinterpret_cast<uint16_t &>(Val) & 0x7f80) >> 7;
+
+  if (Exp16 == 0x7f) {
+    if (Sign)
+      return append(Buf, "-inf");
+    return append(Buf, "inf");
+  }
+  return 0;
+}
+
 template <typename T>
 EnableIfFP<T, unsigned> floatingPointToDecStr(T AbsVal, char *Digits,
                                               int Precision, bool IsSci) {
@@ -1053,6 +1074,8 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(stream) stream
   friend const stream &operator<<(const stream &, const float &);
   friend const stream &operator<<(const stream &, const double &);
   friend const stream &operator<<(const stream &, const half &);
+  friend const stream &operator<<(const stream &,
+                                  const ext::oneapi::bfloat16 &);
 
   friend const stream &operator<<(const stream &, const stream_manipulator);
 
@@ -1159,6 +1182,14 @@ inline const stream &operator<<(const stream &Out, const half &RHS) {
   return Out;
 }
 
+inline const stream &operator<<(const stream &Out,
+                                const ext::oneapi::bfloat16 &RHS) {
+  detail::writeFloatingPoint<ext::oneapi::bfloat16>(
+      Out.GlobalFlushBuf, Out.FlushBufferSize, Out.WIOffset, Out.get_flags(),
+      Out.get_width(), Out.get_precision(), RHS);
+  return Out;
+}
+
 // Pointer
 
 template <typename ElementType, access::address_space Space,
diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp
index b9ec19748e05..819974d1e83b 100644
--- a/sycl/include/sycl/types.hpp
+++ b/sycl/include/sycl/types.hpp
@@ -75,6 +75,8 @@
 #include <sycl/marray.hpp>                     // for __SYCL_BINOP, __SYCL_...
 #include <sycl/multi_ptr.hpp>                  // for multi_ptr
 
+#include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16
+
 #include <array>       // for array
 #include <assert.h>    // for assert
 #include <cstddef>     // for size_t, NULL, byte
@@ -312,6 +314,9 @@ template <typename Type, int NumElements> class vec {
       std::is_same_v<sycl::detail::half_impl::StorageT,
                      sycl::detail::host_half_impl::half>;
 
+  static constexpr bool IsBfloat16 = 
+      std::is_same_v<DataT, sycl::ext::oneapi::bfloat16>;
+
 #if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
 
   static constexpr size_t AdjustedNum = (NumElements == 3) ? 4 : NumElements;
@@ -495,7 +500,13 @@ template <typename Type, int NumElements> class vec {
   template <size_t... Is>
   constexpr vec(const std::array<vec_data_t<DataT>, NumElements> &Arr,
                 std::index_sequence<Is...>)
-      : m_Data{vec_data_t<DataT>(static_cast<DataT>(Arr[Is]))...} {}
+      : m_Data{([&](vec_data_t<DataT> v) constexpr {
+          if constexpr (std::is_same_v<sycl::ext::oneapi::bfloat16, DataT>)
+            return v.value;
+          // return sycl::ext::oneapi::detail::bfloat16ToBits(v);
+          else
+            return vec_data_t<DataT>(static_cast<DataT>(v));
+        })(Arr[Is])...} {}
 
 public:
   using element_type = DataT;
@@ -580,7 +591,7 @@ template <typename Type, int NumElements> class vec {
   template <typename Ty = DataT>
   typename std::enable_if_t<
       std::is_fundamental_v<vec_data_t<Ty>> ||
-          std::is_same_v<typename std::remove_const_t<Ty>, half>,
+          detail::is_half_or_bf16_v<typename std::remove_const_t<Ty>>,
       vec &>
   operator=(const EnableIfNotUsingArrayOnDevice<Ty> &Rhs) {
     m_Data = (DataType)vec_data<Ty>::get(Rhs);
@@ -596,7 +607,7 @@ template <typename Type, int NumElements> class vec {
   template <typename Ty = DataT>
   typename std::enable_if_t<
       std::is_fundamental_v<vec_data_t<Ty>> ||
-          std::is_same_v<typename std::remove_const_t<Ty>, half>,
+          detail::is_half_or_bf16_v<typename std::remove_const_t<Ty>>,
       vec &>
   operator=(const EnableIfUsingArrayOnDevice<Ty> &Rhs) {
     for (int i = 0; i < NumElements; ++i) {
@@ -613,7 +624,7 @@ template <typename Type, int NumElements> class vec {
   template <typename Ty = DataT>
   typename std::enable_if_t<
       std::is_fundamental_v<vec_data_t<Ty>> ||
-          std::is_same_v<typename std::remove_const_t<Ty>, half>,
+          detail::is_half_or_bf16_v<typename std::remove_const_t<Ty>>,
       vec &>
   operator=(const DataT &Rhs) {
     for (int i = 0; i < NumElements; ++i) {
@@ -962,7 +973,7 @@ template <typename Type, int NumElements> class vec {
   typename std::enable_if_t<                                                   \
       std::is_convertible_v<DataT, T> &&                                       \
           (std::is_fundamental_v<vec_data_t<T>> ||                             \
-           std::is_same_v<typename std::remove_const_t<T>, half>),             \
+           detail::is_half_or_bf16_v<typename std::remove_const_t<T>>),        \
       vec>                                                                     \
   operator BINOP(const T & Rhs) const {                                        \
     return *this BINOP vec(static_cast<const DataT &>(Rhs));                   \
@@ -995,7 +1006,7 @@ template <typename Type, int NumElements> class vec {
   typename std::enable_if_t<                                                   \
       std::is_convertible_v<DataT, T> &&                                       \
           (std::is_fundamental_v<vec_data_t<T>> ||                             \
-           std::is_same_v<typename std::remove_const_t<T>, half>),             \
+           detail::is_half_or_bf16_v<typename std::remove_const_t<T>>),        \
       vec>                                                                     \
   operator BINOP(const T & Rhs) const {                                        \
     return *this BINOP vec(static_cast<const DataT &>(Rhs));                   \
@@ -1098,7 +1109,7 @@ template <typename Type, int NumElements> class vec {
   template <typename T>                                                        \
   typename std::enable_if_t<std::is_convertible_v<T, DataT> &&                 \
                                 (std::is_fundamental_v<vec_data_t<T>> ||       \
-                                 std::is_same_v<T, half>),                     \
+                                 detail::is_half_or_bf16_v<T>),                \
                             vec<rel_t, NumElements>>                           \
   operator RELLOGOP(const T & Rhs) const {                                     \
     return *this RELLOGOP vec(static_cast<const DataT &>(Rhs));                \
@@ -1116,7 +1127,7 @@ template <typename Type, int NumElements> class vec {
   template <typename T>                                                        \
   typename std::enable_if_t<std::is_convertible_v<T, DataT> &&                 \
                                 (std::is_fundamental_v<vec_data_t<T>> ||       \
-                                 std::is_same_v<T, half>),                     \
+                                 detail::is_half_or_bf16_v<T>),                \
                             vec<rel_t, NumElements>>                           \
   operator RELLOGOP(const T & Rhs) const {                                     \
     return *this RELLOGOP vec(static_cast<const DataT &>(Rhs));                \
@@ -1258,17 +1269,48 @@ template <typename Type, int NumElements> class vec {
 
   // operator -
   template <typename T = vec> EnableIfNotUsingArray<T> operator-() const {
-    vec Ret{-m_Data};
-    if constexpr (std::is_same_v<Type, bool>) {
-      Ret.ConvertToDataT();
+    namespace oneapi = sycl::ext::oneapi;
+    if constexpr (IsBfloat16 && NumElements == 1) {
+      vec Ret{};
+      oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data);
+      oneapi::bfloat16 w = -v;
+      Ret.m_Data = oneapi::detail::bfloat16ToBits(w);
+    } else if constexpr (IsBfloat16) {
+      vec Ret{};
+      for (size_t I = 0; I < NumElements; ++I) {
+        oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data[I]);
+        oneapi::bfloat16 w = -v;
+        Ret.m_Data[I] = oneapi::detail::bfloat16ToBits(w);
+      }
+      return Ret;
+    }
+    else {
+      vec Ret{-m_Data};
+      if constexpr (std::is_same_v<Type, bool>) {
+        Ret.ConvertToDataT();
+      }
+      return Ret;
     }
-    return Ret;
   }
 
   template <typename T = vec> EnableIfUsingArray<T> operator-() const {
+    namespace oneapi = sycl::ext::oneapi;
     vec Ret{};
-    for (size_t I = 0; I < NumElements; ++I)
-      Ret.setValue(I, vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
+    if constexpr (IsBfloat16 && NumElements == 1) {
+      oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data);
+      oneapi::bfloat16 w = -v;
+      Ret.m_Data = oneapi::detail::bfloat16ToBits(w);
+    } else if constexpr (IsBfloat16) {
+      for (size_t I = 0; I < NumElements; I++) {
+        oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data[I]);
+        oneapi::bfloat16 w = -v;
+        Ret.m_Data[I] = oneapi::detail::bfloat16ToBits(w);
+      }
+    } else {
+      for (size_t I = 0; I < NumElements; ++I)
+        Ret.setValue(I,
+                     vec_data<DataT>::get(-vec_data<DataT>::get(getValue(I))));
+    }
     return Ret;
   }
 #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
@@ -1376,8 +1418,6 @@ template <typename Type, int NumElements> class vec {
 
 #endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
 
-  // CP ---------------
-
   // OP is: &&, ||
   // vec<RET, NumElements> operatorOP(const vec<DataT, NumElements> &Rhs) const;
   // vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
@@ -1614,13 +1654,13 @@ class SwizzleOp {
   using EnableIfScalarType = typename std::enable_if_t<
       std::is_convertible_v<DataT, T> &&
       (std::is_fundamental_v<vec_data_t<T>> ||
-       std::is_same_v<typename std::remove_const_t<T>, half>)>;
+       detail::is_half_or_bf16_v<typename std::remove_const_t<T>>)>;
 
   template <typename T>
   using EnableIfNoScalarType = typename std::enable_if_t<
       !std::is_convertible_v<DataT, T> ||
       !(std::is_fundamental_v<vec_data_t<T>> ||
-        std::is_same_v<typename std::remove_const_t<T>, half>)>;
+        detail::is_half_or_bf16_v<typename std::remove_const_t<T>>)>;
 
   template <int... Indices>
   using Swizzle =
@@ -2196,7 +2236,7 @@ class SwizzleOp {
   template <typename T, int Num>                                               \
   typename std::enable_if_t<                                                   \
       std::is_fundamental_v<vec_data_t<T>> ||                                  \
-          std::is_same_v<typename std::remove_const_t<T>, half>,               \
+          detail::is_half_or_bf16_v<typename std::remove_const_t<T>>,          \
       vec<T, Num>>                                                             \
   operator BINOP(const T & Lhs, const vec<T, Num> &Rhs) {                      \
     return vec<T, Num>(Lhs) BINOP Rhs;                                         \
@@ -2208,7 +2248,7 @@ class SwizzleOp {
   typename std::enable_if_t<                                                   \
       std::is_convertible_v<T, T1> &&                                          \
           (std::is_fundamental_v<vec_data_t<T>> ||                             \
-           std::is_same_v<typename std::remove_const_t<T>, half>),             \
+           detail::is_half_or_bf16_v<typename std::remove_const_t<T>>),        \
       vec<T1, Num>>                                                            \
   operator BINOP(                                                              \
       const T & Lhs,                                                           \
@@ -2252,7 +2292,7 @@ __SYCL_BINOP(<<)
   typename std::enable_if_t<                                                   \
       std::is_convertible_v<T, DataT> &&                                       \
           (std::is_fundamental_v<vec_data_t<T>> ||                             \
-           std::is_same_v<typename std::remove_const_t<T>, half>),             \
+           detail::is_half_or_bf16_v<typename std::remove_const_t<T>>),        \
       vec<detail::rel_t<DataT>, Num>>                                          \
   operator RELLOGOP(const T & Lhs, const vec<DataT, Num> &Rhs) {               \
     return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs;                      \
@@ -2264,7 +2304,7 @@ __SYCL_BINOP(<<)
   typename std::enable_if_t<                                                   \
       std::is_convertible_v<T, T1> &&                                          \
           (std::is_fundamental_v<vec_data_t<T>> ||                             \
-           std::is_same_v<typename std::remove_const_t<T>, half>),             \
+           detail::is_half_or_bf16_v<typename std::remove_const_t<T>>),        \
       vec<detail::rel_t<T1>, Num>>                                             \
   operator RELLOGOP(                                                           \
       const T & Lhs,                                                           \
@@ -2440,10 +2480,10 @@ struct VecStorage<T, 1, typename std::enable_if_t<is_sugeninteger_v<T>>> {
 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
 };
 
-// Single element floating-point (except half)
+// Single element floating-point (except half/bfloat16)
 template <typename T>
 struct VecStorage<
-    T, 1, typename std::enable_if_t<!is_half_v<T> && is_sgenfloat_v<T>>> {
+    T, 1, typename std::enable_if_t<!is_half_or_bf16_v<T> && is_sgenfloat_v<T>>> {
   using DataType = T;
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
 #ifdef __SYCL_DEVICE_ONLY__
@@ -2451,13 +2491,13 @@ struct VecStorage<
 #endif // __SYCL_DEVICE_ONLY__
 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
 };
-// Multiple elements signed/unsigned integers and floating-point (except half)
+// Multiple elements signed/unsigned integers and floating-point (except half/bfloat16)
 template <typename T, int N>
 struct VecStorage<
     T, N,
     typename std::enable_if_t<isValidVectorSize(N) &&
                               (is_sgeninteger_v<T> ||
-                               (is_sgenfloat_v<T> && !is_half_v<T>))>> {
+                               (is_sgenfloat_v<T> && !is_half_or_bf16_v<T>))>> {
   using DataType =
       typename VecStorageImpl<typename VecStorage<T, 1>::DataType, N>::DataType;
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
@@ -2501,6 +2541,25 @@ __SYCL_DEFINE_HALF_VECSTORAGE(4)
 __SYCL_DEFINE_HALF_VECSTORAGE(8)
 __SYCL_DEFINE_HALF_VECSTORAGE(16)
 #undef __SYCL_DEFINE_HALF_VECSTORAGE
+
+// Single element bfloat16
+template <> struct VecStorage<sycl::ext::oneapi::bfloat16, 1, void> {
+  using DataType = sycl::ext::oneapi::detail::Bfloat16StorageT;
+  using VectorDataType = sycl::ext::oneapi::detail::Bfloat16StorageT;
+};
+// Multiple elements bfloat16
+#define __SYCL_DEFINE_BF16_VECSTORAGE(Num)                                     \
+  template <> struct VecStorage<sycl::ext::oneapi::bfloat16, Num, void> {      \
+    using DataType = sycl::ext::oneapi::detail::bf16::Vec##Num##StorageT;      \
+    using VectorDataType =                                                     \
+        sycl::ext::oneapi::detail::bf16::Vec##Num##StorageT;                   \
+  };
+__SYCL_DEFINE_BF16_VECSTORAGE(2)
+__SYCL_DEFINE_BF16_VECSTORAGE(3)
+__SYCL_DEFINE_BF16_VECSTORAGE(4)
+__SYCL_DEFINE_BF16_VECSTORAGE(8)
+__SYCL_DEFINE_BF16_VECSTORAGE(16)
+#undef __SYCL_DEFINE_BF16_VECSTORAGE
 } // namespace detail
 
 /// This macro must be defined to 1 when SYCL implementation allows user
diff --git a/sycl/test/basic_tests/generic_type_traits.cpp b/sycl/test/basic_tests/generic_type_traits.cpp
index 307dc6c4212c..7805adec4578 100644
--- a/sycl/test/basic_tests/generic_type_traits.cpp
+++ b/sycl/test/basic_tests/generic_type_traits.cpp
@@ -63,6 +63,10 @@ int main() {
 
   static_assert(d::is_half_v<s::half>);
 
+  static_assert(d::is_bfloat16_v<sycl::ext::oneapi::bfloat16>);
+  static_assert(d::is_half_or_bf16_v<s::half>);
+  static_assert(d::is_half_or_bf16_v<sycl::ext::oneapi::bfloat16>);
+
   // TODO add checks for the following type traits
   /*
   is_doublen
diff --git a/sycl/test/basic_tests/implicit_device_copyable_types.cpp b/sycl/test/basic_tests/implicit_device_copyable_types.cpp
index dd15f80b8120..ee368207bc24 100644
--- a/sycl/test/basic_tests/implicit_device_copyable_types.cpp
+++ b/sycl/test/basic_tests/implicit_device_copyable_types.cpp
@@ -100,6 +100,9 @@ int main() {
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
   // Extra checks
   static_assert(sycl::is_device_copyable_v<sycl::vec<int, 4>>);
+  static_assert(sycl::is_device_copyable_v<sycl::vec<sycl::half, 4>>);
+  static_assert(
+      sycl::is_device_copyable_v<sycl::vec<sycl::ext::oneapi::bfloat16, 4>>);
 
   struct S {
     sycl::vec<int, 4> v;
diff --git a/sycl/test/basic_tests/types.cpp b/sycl/test/basic_tests/types.cpp
index cc07202f393d..6aab1e433c7a 100644
--- a/sycl/test/basic_tests/types.cpp
+++ b/sycl/test/basic_tests/types.cpp
@@ -204,6 +204,18 @@ void checkVariousVecUnaryOps() {
   checkVecUnaryOps(vd1);
   sycl::vec<double, 16> vd{0, 4, 5, -9, 0, 4, 5, -9, 0, 4, 5, -9, 0, 4, 5, -9};
   checkVecUnaryOps(vd);
+
+  sycl::vec<sycl::half, 1> vh1{1};
+  checkVecUnaryOps(vh1);
+  sycl::vec<sycl::half, 16> vh{0, 4, 5, -9, 0, 4, 5, -9,
+                               0, 4, 5, -9, 0, 4, 5, -9};
+  checkVecUnaryOps(vh);
+
+  sycl::vec<sycl::ext::oneapi::bfloat16, 1> vbf1{1};
+  checkVecUnaryOps(vbf1);
+  sycl::vec<sycl::ext::oneapi::bfloat16, 16> vbf{0, 4, 5, -9, 0, 4, 5, -9,
+                                                 0, 4, 5, -9, 0, 4, 5, -9};
+  checkVecUnaryOps(vbf);
 }
 
 int main() {