[SYCL] Cherry-pick "Switch to use plain array in sycl::vec in more cases" to sycl-rel-6_0_0

sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp

@@ -0,0 +1,134 @@
+//==---------- Forward declarations and traits for vector/marray types -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <cstddef>
+#include <type_traits>
+
+#include <sycl/detail/defines_elementary.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+template <typename DataT, int NumElements> class __SYCL_EBO vec;
+
+template <typename DataT, std::size_t N> class marray;
+
+namespace detail {
+template <typename VecT, typename OperationLeftT, typename OperationRightT,
+          template <typename> class OperationCurrentT, int... Indexes>
+class SwizzleOp;
+
+// Utility for converting a swizzle to a vector or preserve the type if it isn't
+// a swizzle.
+template <typename T> struct simplify_if_swizzle {
+  using type = T;
+};
+
+template <typename VecT, typename OperationLeftT, typename OperationRightT,
+          template <typename> class OperationCurrentT, int... Indexes>
+struct simplify_if_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
+                                     OperationCurrentT, Indexes...>> {
+  using type = vec<typename VecT::element_type, sizeof...(Indexes)>;
+};
+
+template <typename T>
+using simplify_if_swizzle_t = typename simplify_if_swizzle<T>::type;
+
+// --------- is_* traits ------------------ //
+template <typename> struct is_vec : std::false_type {};
+template <typename T, int N> struct is_vec<vec<T, N>> : std::true_type {};
+template <typename T> constexpr bool is_vec_v = is_vec<T>::value;
+
+template <typename T, typename = void>
+struct is_ext_vector : std::false_type {};
+template <typename T, typename = void>
+struct is_valid_type_for_ext_vector : std::false_type {};
+#if defined(__has_extension)
+#if __has_extension(attribute_ext_vector_type)
+template <typename T, int N>
+using ext_vector = T __attribute__((ext_vector_type(N)));
+template <typename T, int N>
+struct is_ext_vector<ext_vector<T, N>> : std::true_type {};
+template <typename T>
+struct is_valid_type_for_ext_vector<T, std::void_t<ext_vector<T, 2>>>
+    : std::true_type {};
+#endif
+#endif
+template <typename T>
+inline constexpr bool is_ext_vector_v = is_ext_vector<T>::value;
+template <typename T>
+inline constexpr bool is_valid_type_for_ext_vector_v =
+    is_valid_type_for_ext_vector<T>::value;
+
+template <typename> struct is_swizzle : std::false_type {};
+template <typename VecT, typename OperationLeftT, typename OperationRightT,
+          template <typename> class OperationCurrentT, int... Indexes>
+struct is_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
+                            OperationCurrentT, Indexes...>> : std::true_type {};
+template <typename T> constexpr bool is_swizzle_v = is_swizzle<T>::value;
+
+template <typename T>
+constexpr bool is_vec_or_swizzle_v = is_vec_v<T> || is_swizzle_v<T>;
+
+template <typename> struct is_marray : std::false_type {};
+template <typename T, std::size_t N>
+struct is_marray<marray<T, N>> : std::true_type {};
+template <typename T> constexpr bool is_marray_v = is_marray<T>::value;
+
+// --------- num_elements trait ------------------ //
+template <typename T>
+struct num_elements : std::integral_constant<std::size_t, 1> {};
+template <typename T, std::size_t N>
+struct num_elements<marray<T, N>> : std::integral_constant<std::size_t, N> {};
+template <typename T, int N>
+struct num_elements<vec<T, N>>
+    : std::integral_constant<std::size_t, std::size_t(N)> {};
+#if defined(__has_extension)
+#if __has_extension(attribute_ext_vector_type)
+template <typename T, int N>
+struct num_elements<T __attribute__((ext_vector_type(N)))>
+    : std::integral_constant<std::size_t, N> {};
+#endif
+#endif
+template <typename VecT, typename OperationLeftT, typename OperationRightT,
+          template <typename> class OperationCurrentT, int... Indexes>
+struct num_elements<SwizzleOp<VecT, OperationLeftT, OperationRightT,
+                              OperationCurrentT, Indexes...>>
+    : std::integral_constant<std::size_t, sizeof...(Indexes)> {};
+
+template <typename T>
+inline constexpr std::size_t num_elements_v = num_elements<T>::value;
+
+// --------- element_type trait ------------------ //
+template <typename T, typename = void> struct element_type {
+  using type = T;
+};
+template <typename T, int N> struct element_type<vec<T, N>> {
+  using type = T;
+};
+template <typename T, std::size_t N> struct element_type<marray<T, N>> {
+  using type = T;
+};
+#if defined(__has_extension)
+#if __has_extension(attribute_ext_vector_type)
+template <typename T, int N>
+struct element_type<T __attribute__((ext_vector_type(N)))> {
+  using type = T;
+};
+#endif
+#endif
+template <typename T> using element_type_t = typename element_type<T>::type;
+
+template <int N>
+inline constexpr bool is_allowed_vec_size_v =
+    N == 1 || N == 2 || N == 3 || N == 4 || N == 8 || N == 16;
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/vector.hpp

@@ -22,6 +22,15 @@
 #endif
 #endif // __clang__
 
+// See vec::DataType definitions for more details
+#ifndef __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE
+#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
+#define __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE 1
+#else
+#define __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE 0
+#endif
+#endif
+
 #if !defined(__HAS_EXT_VECTOR_TYPE__) && defined(__SYCL_DEVICE_ONLY__)
 #error "SYCL device compiler is built without ext_vector_type support"
 #endif
@@ -84,6 +93,9 @@ struct elem {
 };
 
 namespace detail {
+// To be defined in tests, trick to access vec's private methods
+template <typename T1, int T2> class vec_base_test;
+
 template <typename VecT, typename OperationLeftT, typename OperationRightT,
           template <typename> class OperationCurrentT, int... Indexes>
 class SwizzleOp;
@@ -142,7 +154,34 @@ class __SYCL_EBO vec
 
   // This represent type of underlying value. There should be only one field
   // in the class, so vec<float, 16> should be equal to float16 in memory.
-  using DataType = std::array<DataT, AdjustedNum>;
+  //
+  // In intel/llvm#14130 we incorrectly used std::array as an underlying storage
+  // for vec data. The problem with std::array is that it comes from the C++
+  // STL headers which we do not control and they may use something that is
+  // illegal in SYCL device code. One of specific examples is use of debug
+  // assertions in MSVC's STL implementation.
+  //
+  // The better approach is to use plain C++ array, but the problem here is that
+  // C++ specification does not provide any guarantees about the memory layout
+  // of std::array and therefore directly switching to it would technically be
+  // an ABI-break, even though the practical chances of encountering the issue
+  // are low.
+  //
+  // To play it safe, we only switch to use plain array if both its size and
+  // alignment match those of std::array, or unless the new behavior is forced
+  // via __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE or preview breaking changes mode.
+  using DataType = std::conditional_t<
+#if __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE
+      true,
+#else
+      sizeof(std::array<DataT, AdjustedNum>) == sizeof(DataT[AdjustedNum]) &&
+          alignof(std::array<DataT, AdjustedNum>) ==
+              alignof(DataT[AdjustedNum]),
+#endif
+      DataT[AdjustedNum], std::array<DataT, AdjustedNum>>;
+
+  // To allow testing of private methods
+  template <typename T1, int T2> friend class detail::vec_base_test;
 
 #ifdef __SYCL_DEVICE_ONLY__
   using element_type_for_vector_t = typename detail::map_type<

sycl/test/abi/layout_vec.cpp

@@ -12,8 +12,7 @@ SYCL_EXTERNAL void foo(sycl::vec<int, 4>) {}
 
 // CHECK:      0 | class sycl::vec<int, 4>
 // ignore empty base classes
-// CHECK:      0 |   struct std::array<int, 4> m_Data
-// CHECK-NEXT: 0 |     typename {{.+}}::_Type _M_elems
+// CHECK:      0 |   DataType m_Data
 // CHECK-NEXT:   | [sizeof=16, dsize=16, align=16,
 // CHECK-NEXT:   |  nvsize=16, nvalign=16]
 
@@ -23,7 +22,6 @@ SYCL_EXTERNAL void foo(sycl::vec<bool, 16>) {}
 
 // CHECK:      0 | class sycl::vec<_Bool, 16>
 // ignore empty base classes
-// CHECK:      0 |   struct std::array<_Bool, 16> m_Data
-// CHECK-NEXT: 0 |     typename {{.+}}::_Type _M_elems
+// CHECK:      0 |   DataType m_Data
 // CHECK-NEXT:   | [sizeof=16, dsize=16, align=16,
 // CHECK-NEXT:   |  nvsize=16, nvalign=16]

sycl/test/basic_tests/vectors/storage.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -Xclang -verify %s -fsyntax-only
+// RUN: %clangxx -fsycl -Xclang -verify %s -fsyntax-only -fpreview-breaking-changes
+// RUN: %clangxx -fsycl -Xclang -verify %s -fsyntax-only -D__SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE=1
+// expected-no-diagnostics
+
+#include <sycl/vector.hpp>
+
+#include <type_traits>
+
+namespace sycl {
+namespace detail {
+template <typename T, int N> class vec_base_test {
+public:
+  static void do_check() {
+    constexpr bool uses_std_array =
+        std::is_same_v<typename sycl::vec<T, N>::DataType, std::array<T, N>>;
+    constexpr bool uses_plain_array =
+        std::is_same_v<typename sycl::vec<T, N>::DataType, T[N]>;
+
+    constexpr bool std_array_and_plain_array_have_the_same_layout =
+        sizeof(std::array<T, N>) == sizeof(T[N]) &&
+        alignof(std::array<T, N>) == alignof(T[N]);
+
+#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) ||                               \
+    __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE
+    static_assert(uses_plain_array,
+                  "We must use plain array regardless of "
+                  "layout, because user is opted-in for a potential ABI-break");
+#else
+    static_assert(std_array_and_plain_array_have_the_same_layout ==
+                      uses_plain_array,
+                  "If layouts are the same, we must use safer plain array "
+                  "instead of std::array, or vice versa");
+    static_assert(
+        !std_array_and_plain_array_have_the_same_layout == uses_std_array,
+        "If layouts are not the same, we must use std::array to preserve ABI");
+#endif
+  }
+};
+} // namespace detail
+} // namespace sycl
+
+int main() { sycl::detail::vec_base_test<int, 4>::do_check(); }

sycl/test/regression/vec_array_windows.cpp

@@ -0,0 +1,23 @@
+// Test to isolate sycl::vec regression after
+// https://github.com/intel/llvm/pull/14130. This PR caused sycl::vec to use
+// std::array as its underlying storage. However, operations on std::array
+// may emit debug-mode-only functions, on which the device compiler may fail.
+
+// REQUIRES: windows
+
+// RUN: %clangxx -fsycl -D_DEBUG %s -fsycl-device-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
+// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes -D_DEBUG -fsycl-device-only %s %}
+
+#include <sycl/sycl.hpp>
+
+// expected-no-diagnostics
+//
+// Our current implementation automatically opts-in for a new implementation if
+// that is possible without breaking ABI.
+// However, depending on the environment (used STL implementation, in
+// particular) it may not be the case. Therefore, the lines below are kept for
+// reference of how an error would look like in a problematic environment.
+// not-expected-error@* {{SYCL kernel cannot call a variadic function}}
+// not-expected-error@* {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
+// not-expected-error@* {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
+SYCL_EXTERNAL auto GetFirstElement(sycl::vec<int, 3> v) { return v[0]; }