Skip to content

Commit b7d5141

Browse files
AlexeySachkovAlexeySachkov
committed
[SYCL] Switch to use plain array in sycl::vec in more cases
The problem with using `std::array` in `sycl::vec` is that we cannot compile it in some environments (namely, Windows) because the former may use something that is illegal in SYCL device code. intel#17025 fixed that, but only did so under preview breaking changes mode, which does not satisfy some of our customers immediately. This PR introduces two main changes: - it allows to opt-in for new behavior through passing `-D__SYCL_USE_NEW_VEC_IMPL=1` macro without using `-fpreview-breaking-changes` flag. That allows for a more gradual opt-in from customers who are interested in this fix - it switches the imlpementation to use the new approach with C-style arrays if their size and alignment is the same as for the corresponding `std::array` - in that case their memory layout is expected to be absolutely the same and therefore it should be safe to use the new approach without fear of some ABI incompatibilities. This allows for customers to benefit from the fix without specifying any extra macro (which should be the case for the most common platforms out there) This is a cherry-pick of intel#17656
1 parent 54cbbec commit b7d5141

File tree

5 files changed

+242
-5
lines changed

5 files changed

+242
-5
lines changed

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

Lines changed: 134 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,134 @@
1+
//==---------- Forward declarations and traits for vector/marray types -----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <cstddef>
12+
#include <type_traits>
13+
14+
#include <sycl/detail/defines_elementary.hpp>
15+
16+
namespace sycl {
17+
inline namespace _V1 {
18+
template <typename DataT, int NumElements> class __SYCL_EBO vec;
19+
20+
template <typename DataT, std::size_t N> class marray;
21+
22+
namespace detail {
23+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
24+
template <typename> class OperationCurrentT, int... Indexes>
25+
class SwizzleOp;
26+
27+
// Utility for converting a swizzle to a vector or preserve the type if it isn't
28+
// a swizzle.
29+
template <typename T> struct simplify_if_swizzle {
30+
using type = T;
31+
};
32+
33+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
34+
template <typename> class OperationCurrentT, int... Indexes>
35+
struct simplify_if_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
36+
OperationCurrentT, Indexes...>> {
37+
using type = vec<typename VecT::element_type, sizeof...(Indexes)>;
38+
};
39+
40+
template <typename T>
41+
using simplify_if_swizzle_t = typename simplify_if_swizzle<T>::type;
42+
43+
// --------- is_* traits ------------------ //
44+
template <typename> struct is_vec : std::false_type {};
45+
template <typename T, int N> struct is_vec<vec<T, N>> : std::true_type {};
46+
template <typename T> constexpr bool is_vec_v = is_vec<T>::value;
47+
48+
template <typename T, typename = void>
49+
struct is_ext_vector : std::false_type {};
50+
template <typename T, typename = void>
51+
struct is_valid_type_for_ext_vector : std::false_type {};
52+
#if defined(__has_extension)
53+
#if __has_extension(attribute_ext_vector_type)
54+
template <typename T, int N>
55+
using ext_vector = T __attribute__((ext_vector_type(N)));
56+
template <typename T, int N>
57+
struct is_ext_vector<ext_vector<T, N>> : std::true_type {};
58+
template <typename T>
59+
struct is_valid_type_for_ext_vector<T, std::void_t<ext_vector<T, 2>>>
60+
: std::true_type {};
61+
#endif
62+
#endif
63+
template <typename T>
64+
inline constexpr bool is_ext_vector_v = is_ext_vector<T>::value;
65+
template <typename T>
66+
inline constexpr bool is_valid_type_for_ext_vector_v =
67+
is_valid_type_for_ext_vector<T>::value;
68+
69+
template <typename> struct is_swizzle : std::false_type {};
70+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
71+
template <typename> class OperationCurrentT, int... Indexes>
72+
struct is_swizzle<SwizzleOp<VecT, OperationLeftT, OperationRightT,
73+
OperationCurrentT, Indexes...>> : std::true_type {};
74+
template <typename T> constexpr bool is_swizzle_v = is_swizzle<T>::value;
75+
76+
template <typename T>
77+
constexpr bool is_vec_or_swizzle_v = is_vec_v<T> || is_swizzle_v<T>;
78+
79+
template <typename> struct is_marray : std::false_type {};
80+
template <typename T, std::size_t N>
81+
struct is_marray<marray<T, N>> : std::true_type {};
82+
template <typename T> constexpr bool is_marray_v = is_marray<T>::value;
83+
84+
// --------- num_elements trait ------------------ //
85+
template <typename T>
86+
struct num_elements : std::integral_constant<std::size_t, 1> {};
87+
template <typename T, std::size_t N>
88+
struct num_elements<marray<T, N>> : std::integral_constant<std::size_t, N> {};
89+
template <typename T, int N>
90+
struct num_elements<vec<T, N>>
91+
: std::integral_constant<std::size_t, std::size_t(N)> {};
92+
#if defined(__has_extension)
93+
#if __has_extension(attribute_ext_vector_type)
94+
template <typename T, int N>
95+
struct num_elements<T __attribute__((ext_vector_type(N)))>
96+
: std::integral_constant<std::size_t, N> {};
97+
#endif
98+
#endif
99+
template <typename VecT, typename OperationLeftT, typename OperationRightT,
100+
template <typename> class OperationCurrentT, int... Indexes>
101+
struct num_elements<SwizzleOp<VecT, OperationLeftT, OperationRightT,
102+
OperationCurrentT, Indexes...>>
103+
: std::integral_constant<std::size_t, sizeof...(Indexes)> {};
104+
105+
template <typename T>
106+
inline constexpr std::size_t num_elements_v = num_elements<T>::value;
107+
108+
// --------- element_type trait ------------------ //
109+
template <typename T, typename = void> struct element_type {
110+
using type = T;
111+
};
112+
template <typename T, int N> struct element_type<vec<T, N>> {
113+
using type = T;
114+
};
115+
template <typename T, std::size_t N> struct element_type<marray<T, N>> {
116+
using type = T;
117+
};
118+
#if defined(__has_extension)
119+
#if __has_extension(attribute_ext_vector_type)
120+
template <typename T, int N>
121+
struct element_type<T __attribute__((ext_vector_type(N)))> {
122+
using type = T;
123+
};
124+
#endif
125+
#endif
126+
template <typename T> using element_type_t = typename element_type<T>::type;
127+
128+
template <int N>
129+
inline constexpr bool is_allowed_vec_size_v =
130+
N == 1 || N == 2 || N == 3 || N == 4 || N == 8 || N == 16;
131+
132+
} // namespace detail
133+
} // namespace _V1
134+
} // namespace sycl

‎sycl/include/sycl/vector.hpp

Lines changed: 40 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,15 @@
2222
#endif
2323
#endif // __clang__
2424

25+
// See vec::DataType definitions for more details
26+
#ifndef __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE
27+
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
28+
#define __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE 1
29+
#else
30+
#define __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE 0
31+
#endif
32+
#endif
33+
2534
#if !defined(__HAS_EXT_VECTOR_TYPE__) && defined(__SYCL_DEVICE_ONLY__)
2635
#error "SYCL device compiler is built without ext_vector_type support"
2736
#endif
@@ -84,6 +93,9 @@ struct elem {
8493
};
8594

8695
namespace detail {
96+
// To be defined in tests, trick to access vec's private methods
97+
template <typename T1, int T2> class vec_base_test;
98+
8799
template <typename VecT, typename OperationLeftT, typename OperationRightT,
88100
template <typename> class OperationCurrentT, int... Indexes>
89101
class SwizzleOp;
@@ -142,7 +154,34 @@ class __SYCL_EBO vec
142154

143155
// This represent type of underlying value. There should be only one field
144156
// in the class, so vec<float, 16> should be equal to float16 in memory.
145-
using DataType = std::array<DataT, AdjustedNum>;
157+
//
158+
// In intel/llvm#14130 we incorrectly used std::array as an underlying storage
159+
// for vec data. The problem with std::array is that it comes from the C++
160+
// STL headers which we do not control and they may use something that is
161+
// illegal in SYCL device code. One of specific examples is use of debug
162+
// assertions in MSVC's STL implementation.
163+
//
164+
// The better approach is to use plain C++ array, but the problem here is that
165+
// C++ specification does not provide any guarantees about the memory layout
166+
// of std::array and therefore directly switching to it would technically be
167+
// an ABI-break, even though the practical chances of encountering the issue
168+
// are low.
169+
//
170+
// To play it safe, we only switch to use plain array if both its size and
171+
// alignment match those of std::array, or unless the new behavior is forced
172+
// via __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE or preview breaking changes mode.
173+
using DataType = std::conditional_t<
174+
#if __SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE
175+
true,
176+
#else
177+
sizeof(std::array<DataT, AdjustedNum>) == sizeof(DataT[AdjustedNum]) &&
178+
alignof(std::array<DataT, AdjustedNum>) ==
179+
alignof(DataT[AdjustedNum]),
180+
#endif
181+
DataT[AdjustedNum], std::array<DataT, AdjustedNum>>;
182+
183+
// To allow testing of private methods
184+
template <typename T1, int T2> friend class detail::vec_base_test;
146185

147186
#ifdef __SYCL_DEVICE_ONLY__
148187
using element_type_for_vector_t = typename detail::map_type<

‎sycl/test/abi/layout_vec.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,7 @@ SYCL_EXTERNAL void foo(sycl::vec<int, 4>) {}
1212

1313
// CHECK: 0 | class sycl::vec<int, 4>
1414
// ignore empty base classes
15-
// CHECK: 0 | struct std::array<int, 4> m_Data
16-
// CHECK-NEXT: 0 | typename {{.+}}::_Type _M_elems
15+
// CHECK: 0 | DataType m_Data
1716
// CHECK-NEXT: | [sizeof=16, dsize=16, align=16,
1817
// CHECK-NEXT: | nvsize=16, nvalign=16]
1918

@@ -23,7 +22,6 @@ SYCL_EXTERNAL void foo(sycl::vec<bool, 16>) {}
2322

2423
// CHECK: 0 | class sycl::vec<_Bool, 16>
2524
// ignore empty base classes
26-
// CHECK: 0 | struct std::array<_Bool, 16> m_Data
27-
// CHECK-NEXT: 0 | typename {{.+}}::_Type _M_elems
25+
// CHECK: 0 | DataType m_Data
2826
// CHECK-NEXT: | [sizeof=16, dsize=16, align=16,
2927
// CHECK-NEXT: | nvsize=16, nvalign=16]

‎sycl/test/basic_tests/vectors/storage.cpp

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
// RUN: %clangxx -fsycl -Xclang -verify %s -fsyntax-only
2+
// RUN: %clangxx -fsycl -Xclang -verify %s -fsyntax-only -fpreview-breaking-changes
3+
// RUN: %clangxx -fsycl -Xclang -verify %s -fsyntax-only -D__SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE=1
4+
// expected-no-diagnostics
5+
6+
#include <sycl/vector.hpp>
7+
8+
#include <type_traits>
9+
10+
namespace sycl {
11+
namespace detail {
12+
template <typename T, int N> class vec_base_test {
13+
public:
14+
static void do_check() {
15+
constexpr bool uses_std_array =
16+
std::is_same_v<typename sycl::vec<T, N>::DataType, std::array<T, N>>;
17+
constexpr bool uses_plain_array =
18+
std::is_same_v<typename sycl::vec<T, N>::DataType, T[N]>;
19+
20+
constexpr bool std_array_and_plain_array_have_the_same_layout =
21+
sizeof(std::array<T, N>) == sizeof(T[N]) &&
22+
alignof(std::array<T, N>) == alignof(T[N]);
23+
24+
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) || \
25+
__SYCL_USE_PLAIN_ARRAY_AS_VEC_STORAGE
26+
static_assert(uses_plain_array,
27+
"We must use plain array regardless of "
28+
"layout, because user is opted-in for a potential ABI-break");
29+
#else
30+
static_assert(std_array_and_plain_array_have_the_same_layout ==
31+
uses_plain_array,
32+
"If layouts are the same, we must use safer plain array "
33+
"instead of std::array, or vice versa");
34+
static_assert(
35+
!std_array_and_plain_array_have_the_same_layout == uses_std_array,
36+
"If layouts are not the same, we must use std::array to preserve ABI");
37+
#endif
38+
}
39+
};
40+
} // namespace detail
41+
} // namespace sycl
42+
43+
int main() { sycl::detail::vec_base_test<int, 4>::do_check(); }

‎sycl/test/regression/vec_array_windows.cpp

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

0 commit comments

Comments
 (0)