Fix failures

intel · Apr 11, 2024 · a75a93b · a75a93b
1 parent 1fe41c2
commit a75a93b
Show file tree

Hide file tree

Showing 3 changed files with 38 additions and 20 deletions.
diff --git a/sycl/include/sycl/ext/oneapi/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/bfloat16.hpp
@@ -38,6 +38,12 @@ 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)));
+
+using Vec2DeviceStorageT = std::array<Bfloat16StorageT, 2>;
+using Vec3DeviceStorageT = std::array<Bfloat16StorageT, 3>;
+using Vec4DeviceStorageT = std::array<Bfloat16StorageT, 4>;
+using Vec8DeviceStorageT = std::array<Bfloat16StorageT, 8>;
+using Vec16DeviceStorageT = std::array<Bfloat16StorageT, 16>;
 #else
 using Vec2StorageT = std::array<Bfloat16StorageT, 2>;
 using Vec3StorageT = std::array<Bfloat16StorageT, 3>;

diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp
@@ -1032,25 +1032,34 @@ template <typename Type, int NumElements> class vec {
   // operator!
   friend vec<detail::rel_t<DataT>, NumElements> operator!(const vec &Rhs) {
     vec Ret{};
-#ifndef __SYCL_DEVICE_ONLY__
-    for (size_t I = 0; I < NumElements; ++I) {
+#ifdef __SYCL_DEVICE_ONLY__
+    if constexpr (!std::is_same_v<DataT, sycl::ext::oneapi::bfloat16> &&
+                  !std::is_same_v<DataT, sycl::detail::half_impl::half> &&
 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-      // std::byte neither supports ! unary op or casting, so special handling
-      // is needed. And, worse, Windows has a conflict with 'byte'.
-      if constexpr (std::is_same_v<std::byte, DataT>) {
-        Ret.setValue(I, std::byte{!vec_data<DataT>::get(Rhs.getValue(I))});
-      } else
-#endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
-      {
-        Ret.setValue(I, !vec_data<DataT>::get(Rhs.getValue(I)));
-      }
-    }
-    return Ret.template as<vec<detail::rel_t<DataT>, NumElements>>();
+                  !std::is_same_v<DataT, std::byte>)   
 #else
+                  1)
+#endif         
+  {
     auto extVec = detail::BitCast<DataType, vector_t>(Rhs.m_Data);
-    return vec{!extVec}
-        .template as<vec<detail::rel_t<DataT>, NumElements>>();
+    return vec<detail::rel_t<DataT>, NumElements>{!extVec};
+  } else
 #endif
+    {
+      for (size_t I = 0; I < NumElements; ++I) {
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+        // std::byte neither supports ! unary op or casting, so special handling
+        // is needed. And, worse, Windows has a conflict with 'byte'.
+        if constexpr (std::is_same_v<std::byte, DataT>) {
+          Ret.setValue(I, std::byte{!vec_data<DataT>::get(Rhs.getValue(I))});
+        } else
+#endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+        {
+          Ret.setValue(I, !vec_data<DataT>::get(Rhs.getValue(I)));
+        }
+      }
+      return Ret.template as<vec<detail::rel_t<DataT>, NumElements>>();
+    }
   }
 
   // operator +
@@ -2088,12 +2097,18 @@ template <> struct VecStorage<sycl::ext::oneapi::bfloat16, 1, void> {
   using VectorDataType = sycl::ext::oneapi::detail::Bfloat16StorageT;
 };
 // Multiple elements bfloat16
+#ifndef __SYCL_DEVICE_ONLY__
 #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;                   \
   };
+#else
+#define __SYCL_DEFINE_BF16_VECSTORAGE(Num)                                     \
+  template <> struct VecStorage<sycl::ext::oneapi::bfloat16, Num, void> {      \
+    using DataType = sycl::ext::oneapi::detail::bf16::Vec##Num##DeviceStorageT; \
+    using VectorDataType = sycl::ext::oneapi::detail::bf16::Vec##Num##StorageT;  \
+  };
+#endif
 __SYCL_DEFINE_BF16_VECSTORAGE(2)
 __SYCL_DEFINE_BF16_VECSTORAGE(3)
 __SYCL_DEFINE_BF16_VECSTORAGE(4)

diff --git a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order_relational.cpp b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order_relational.cpp
@@ -2,9 +2,6 @@
 // RUN: %{build} -fpreview-breaking-changes -o %t.out
 // RUN: %{run} %t.out
 
-// This test currently fails on AMD HIP due to an unresolved memcmp function.
-// XFAIL: hip_amd
-
 // Checks scalar/vec relational operator ordering.
 
 #include "vec_binary_scalar_order.hpp"