Skip to content

Commit d92787a

Browse files
KornevNikitaAlexeySachkov
KornevNikita
and
AlexeySachkov
authored
[SYCL] Switch to use plain array in sycl::vec in more cases (#17804)
### Cherry-pick of: #17656 and some changes to resolve merge conflicts & make it work. 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. #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) Co-authored-by: Alexey Sachkov <alexey.sachkov@intel.com>
1 parent 57375e4 commit d92787a

File tree

8 files changed

+165
-61
lines changed

8 files changed

+165
-61
lines changed

‎sycl/include/sycl/detail/vector_convert.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -57,6 +57,7 @@
5757
#include <sycl/detail/generic_type_traits.hpp> // for is_sigeninteger, is_s...
5858
#include <sycl/exception.hpp> // for errc
5959

60+
#include <sycl/detail/memcpy.hpp>
6061
#include <sycl/ext/oneapi/bfloat16.hpp> // bfloat16
6162
#include <sycl/vector.hpp>
6263

@@ -941,7 +942,7 @@ vec<convertT, NumElements> vec<DataT, NumElements>::convert() const {
941942
if constexpr (canUseNativeVectorConvert) {
942943
auto val = detail::convertImpl<T, R, roundingMode, NumElements,
943944
OpenCLVecT, OpenCLVecR>(NativeVector);
944-
Result.m_Data = sycl::bit_cast<decltype(Result.m_Data)>(val);
945+
sycl::detail::memcpy_no_adl(&Result.m_Data, &val, sizeof(Result));
945946
} else
946947
#endif // __SYCL_DEVICE_ONLY__
947948
{

‎sycl/include/sycl/vector.hpp

Lines changed: 41 additions & 2 deletions
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
@@ -85,6 +94,9 @@ struct elem {
8594
};
8695

8796
namespace detail {
97+
// To be defined in tests, trick to access vec's private methods
98+
template <typename T1, int T2> class vec_base_test;
99+
88100
template <typename VecT, typename OperationLeftT, typename OperationRightT,
89101
template <typename> class OperationCurrentT, int... Indexes>
90102
class SwizzleOp;
@@ -149,7 +161,34 @@ class __SYCL_EBO vec
149161

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

154193
#ifdef __SYCL_DEVICE_ONLY__
155194
using element_type_for_vector_t = typename detail::map_type<
@@ -318,7 +357,7 @@ class __SYCL_EBO vec
318357
typename vector_t_ = vector_t,
319358
typename = typename std::enable_if_t<std::is_same_v<vector_t_, vector_t>>>
320359
constexpr vec(vector_t_ openclVector) {
321-
m_Data = sycl::bit_cast<DataType>(openclVector);
360+
sycl::detail::memcpy_no_adl(&m_Data, &openclVector, sizeof(openclVector));
322361
}
323362

324363
/* @SYCL2020

‎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 _AT_Type::_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 _AT_Type::_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/check_device_code/vector/vector_bf16_builtins.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ SYCL_EXTERNAL auto TestFMin(vec<bfloat16, 2> a, vec<bfloat16, 2> b) {
6969
}
7070

7171
// CHECK-LABEL: define dso_local spir_func void @_Z8TestFMaxN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEES5_(
72-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.5") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.5") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.5") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] {
72+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.4") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.4") align 8 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.4") align 8 [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] {
7373
// CHECK-NEXT: entry:
7474
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I12_I:%.*]] = alloca <3 x float>, align 16
7575
// CHECK-NEXT: [[DST_I_I_I_I13_I:%.*]] = alloca [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
@@ -123,7 +123,7 @@ SYCL_EXTERNAL auto TestFMax(vec<bfloat16, 3> a, vec<bfloat16, 3> b) {
123123
}
124124

125125
// CHECK-LABEL: define dso_local spir_func void @_Z9TestIsNanN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi4EEE(
126-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.15") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.20") align 8 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META38:![0-9]+]] !sycl_fixed_targets [[META7]] {
126+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.12") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.16") align 8 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META38:![0-9]+]] !sycl_fixed_targets [[META7]] {
127127
// CHECK-NEXT: entry:
128128
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I_I:%.*]] = alloca <4 x i16>, align 8
129129
// CHECK-NEXT: [[DST_I_I_I_I_I:%.*]] = alloca [4 x float], align 4
@@ -149,7 +149,7 @@ SYCL_EXTERNAL auto TestIsNan(vec<bfloat16, 4> a) {
149149
}
150150

151151
// CHECK-LABEL: define dso_local spir_func void @_Z8TestFabsN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE(
152-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.38") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.38") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] {
152+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.32") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.32") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] {
153153
// CHECK-NEXT: entry:
154154
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32
155155
// CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
@@ -185,7 +185,7 @@ SYCL_EXTERNAL auto TestFabs(vec<bfloat16, 8> a) {
185185
}
186186

187187
// CHECK-LABEL: define dso_local spir_func void @_Z8TestCeilN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE(
188-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.38") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.38") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META59:![0-9]+]] !sycl_fixed_targets [[META7]] {
188+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.32") align 16 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.32") align 16 [[A:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META59:![0-9]+]] !sycl_fixed_targets [[META7]] {
189189
// CHECK-NEXT: entry:
190190
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32
191191
// CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
@@ -221,7 +221,7 @@ SYCL_EXTERNAL auto TestCeil(vec<bfloat16, 8> a) {
221221
}
222222

223223
// CHECK-LABEL: define dso_local spir_func void @_Z7TestFMAN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi16EEES5_S5_(
224-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.48") align 32 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.48") align 32 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.48") align 32 [[B:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.48") align 32 [[C:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META70:![0-9]+]] !sycl_fixed_targets [[META7]] {
224+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.40") align 32 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.40") align 32 [[A:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.40") align 32 [[B:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::vec.40") align 32 [[C:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META70:![0-9]+]] !sycl_fixed_targets [[META7]] {
225225
// CHECK-NEXT: entry:
226226
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I14_I:%.*]] = alloca <16 x float>, align 64
227227
// CHECK-NEXT: [[DST_I_I_I_I15_I:%.*]] = alloca [16 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2

‎sycl/test/check_device_code/vector/vector_convert_bfloat.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ SYCL_EXTERNAL auto TestBFtoFDeviceRZ(vec<bfloat16, 3> &inp) {
6363
}
6464

6565
// CHECK-LABEL: define dso_local spir_func void @_Z19TestBFtointDeviceRZRN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEE(
66-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.5") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META18:![0-9]+]] !sycl_fixed_targets [[META7]] {
66+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.4") align 16 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META18:![0-9]+]] !sycl_fixed_targets [[META7]] {
6767
// CHECK-NEXT: entry:
6868
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META19:![0-9]+]])
6969
// CHECK-NEXT: [[LOADVEC4_I_I:%.*]] = load <4 x i16>, ptr addrspace(4) [[INP]], align 8, !noalias [[META19]]
@@ -93,7 +93,7 @@ SYCL_EXTERNAL auto TestBFtointDeviceRZ(vec<bfloat16, 3> &inp) {
9393
}
9494

9595
// CHECK-LABEL: define dso_local spir_func void @_Z20TestBFtointDeviceRNERN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi1EEE(
96-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.10") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 2 dereferenceable(2) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] {
96+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.8") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 2 dereferenceable(2) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META24:![0-9]+]] !sycl_fixed_targets [[META7]] {
9797
// CHECK-NEXT: entry:
9898
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META25:![0-9]+]])
9999
// CHECK-NEXT: [[TMP0:%.*]] = load i16, ptr addrspace(4) [[INP]], align 2, !tbaa [[TBAA11]], !noalias [[META25]]
@@ -191,7 +191,7 @@ SYCL_EXTERNAL auto TestInttoBFDeviceRZ(vec<int, 3> &inp) {
191191
}
192192

193193
// CHECK-LABEL: define dso_local spir_func void @_Z19TestLLtoBFDeviceRTPRN4sycl3_V13vecIxLi1EEE(
194-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.15") align 2 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] {
194+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.12") align 2 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 8 dereferenceable(8) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META42:![0-9]+]] !sycl_fixed_targets [[META7]] {
195195
// CHECK-NEXT: entry:
196196
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META43:![0-9]+]])
197197
// CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(4) [[INP]], align 8, !tbaa [[TBAA46:![0-9]+]], !noalias [[META43]]
@@ -204,7 +204,7 @@ SYCL_EXTERNAL auto TestLLtoBFDeviceRTP(vec<long long, 1> &inp) {
204204
}
205205

206206
// CHECK-LABEL: define dso_local spir_func void @_Z22TestShorttoBFDeviceRTNRN4sycl3_V13vecIsLi2EEE(
207-
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.25") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 4 dereferenceable(4) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] {
207+
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::vec.20") align 4 [[AGG_RESULT:%.*]], ptr addrspace(4) nocapture noundef readonly align 4 dereferenceable(4) [[INP:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META48:![0-9]+]] !sycl_fixed_targets [[META7]] {
208208
// CHECK-NEXT: entry:
209209
// CHECK-NEXT: tail call void @llvm.experimental.noalias.scope.decl(metadata [[META49:![0-9]+]])
210210
// CHECK-NEXT: [[TMP0:%.*]] = load <2 x i16>, ptr addrspace(4) [[INP]], align 4, !tbaa [[TBAA11]], !noalias [[META49]]

0 commit comments

Comments
 (0)