Skip to content

Commit a9d0e1b

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Fix strict alias violations in vec representation (#11883)
Parts of the vector implementation violates strict aliasing by representing vectors of certain types by another element type and reinterpret-casting to the right type to correctly pass references along. This may especially affect fundamental integral types that are not the default fixed-width integer type of the system. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 08907a3 commit a9d0e1b

File tree

1 file changed

+66
-49
lines changed

1 file changed

+66
-49
lines changed

‎sycl/include/sycl/types.hpp

Lines changed: 66 additions & 49 deletions
Original file line numberDiff line numberDiff line change
@@ -314,19 +314,6 @@ template <typename Type, int NumElements> class vec {
314314

315315
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
316316

317-
// This represents HOW we will approach the underlying value, so as to
318-
// benefit from vector speed improvements
319-
using VectorDataType =
320-
typename detail::VecStorage<DataT, NumElements>::VectorDataType;
321-
322-
VectorDataType &getAsVector() {
323-
return *reinterpret_cast<VectorDataType *>(m_Data.data());
324-
}
325-
326-
const VectorDataType &getAsVector() const {
327-
return *reinterpret_cast<const VectorDataType *>(m_Data.data());
328-
}
329-
330317
static constexpr size_t AdjustedNum = (NumElements == 3) ? 4 : NumElements;
331318
static constexpr size_t Sz = sizeof(DataT) * AdjustedNum;
332319
static constexpr bool IsSizeGreaterThanMaxAlign =
@@ -516,7 +503,8 @@ template <typename Type, int NumElements> class vec {
516503

517504
#ifdef __SYCL_DEVICE_ONLY__
518505
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
519-
using vector_t = VectorDataType;
506+
using vector_t =
507+
typename detail::VecStorage<DataT, NumElements>::VectorDataType;
520508
#else
521509
using vector_t = DataType;
522510
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
@@ -713,7 +701,7 @@ template <typename Type, int NumElements> class vec {
713701
if constexpr (!IsUsingArrayOnDevice) {
714702
return m_Data;
715703
} else {
716-
auto ptr = bit_cast<const VectorDataType *>((&m_Data)->data());
704+
auto ptr = bit_cast<const vector_t *>((&m_Data)->data());
717705
return *ptr;
718706
}
719707
}
@@ -788,32 +776,32 @@ template <typename Type, int NumElements> class vec {
788776
"Unsupported convertT");
789777
using T = vec_data_t<DataT>;
790778
using R = vec_data_t<convertT>;
779+
using OpenCLT = detail::ConvertToOpenCLType_t<T>;
780+
using OpenCLR = detail::ConvertToOpenCLType_t<R>;
791781
vec<convertT, NumElements> Result;
792-
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
793-
if constexpr (NativeVec && vec<convertT, NumElements>::NativeVec) {
794-
#ifdef __SYCL_DEVICE_ONLY__
795-
// If both vectors are representable as native vectors, then we can use
796-
// a single vector-wide operation to do a conversion:
797-
Result.m_Data = detail::convertImpl<
798-
T, R, roundingMode, NumElements, VectorDataType,
799-
typename vec<convertT, NumElements>::VectorDataType>(m_Data);
800-
#endif // __SYCL_DEVICE_ONLY
801-
} else {
802-
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
803-
782+
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
783+
using OpenCLVecT = OpenCLT __attribute__((ext_vector_type(NumElements)));
784+
using OpenCLVecR = OpenCLR __attribute__((ext_vector_type(NumElements)));
785+
if constexpr (NativeVec && vec<convertT, NumElements>::NativeVec &&
786+
std::is_convertible_v<decltype(m_Data), OpenCLVecT> &&
787+
std::is_convertible_v<decltype(Result.m_Data), OpenCLR>) {
788+
// If both vectors are representable as native vectors and these native
789+
// vectors can be converted to valid OpenCL representations, then we can
790+
// use a single vector-wide operation to do a conversion:
791+
Result.m_Data = detail::convertImpl<T, R, roundingMode, NumElements,
792+
OpenCLVecT, OpenCLVecR>(m_Data);
793+
} else
794+
#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) &&
795+
// defined(__SYCL_DEVICE_ONLY__)
796+
{
804797
// Otherwise, we fallback to per-element conversion:
805-
using OpenCLT = detail::ConvertToOpenCLType_t<vec_data_t<DataT>>;
806-
using OpenCLR = detail::ConvertToOpenCLType_t<vec_data_t<convertT>>;
807798
for (size_t I = 0; I < NumElements; ++I) {
808799
Result.setValue(
809800
I, vec_data<convertT>::get(
810801
detail::convertImpl<T, R, roundingMode, 1, OpenCLT, OpenCLR>(
811802
vec_data<DataT>::get(getValue(I)))));
812803
}
813-
814-
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
815804
}
816-
#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
817805

818806
if constexpr (std::is_same_v<convertT, bool>) {
819807
Result.ConvertToDataT();
@@ -982,7 +970,7 @@ template <typename Type, int NumElements> class vec {
982970
vec operator BINOP(const vec &Rhs) const { \
983971
vec Ret{}; \
984972
if constexpr (NativeVec) \
985-
Ret.getAsVector() = getAsVector() BINOP Rhs.getAsVector(); \
973+
Ret.m_Data = m_Data BINOP Rhs.m_Data; \
986974
else \
987975
for (size_t I = 0; I < NumElements; ++I) \
988976
Ret.setValue(I, (DataT)(vec_data<DataT>::get(getValue( \
@@ -2258,14 +2246,10 @@ template <typename T, int N> struct VecStorageImpl {
22582246
using VectorDataType = T __attribute__((ext_vector_type(N)));
22592247
};
22602248
#else // __SYCL_DEVICE_ONLY__
2261-
2262-
template <typename T, int N> struct VecStorageImpl;
2263-
#define __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, num) \
2264-
template <> struct VecStorageImpl<type, num> { \
2265-
using DataType = std::array<type, (num == 3) ? 4 : num>; \
2266-
using VectorDataType = ::cl_##cl_type##num; \
2267-
};
2268-
#endif // SYCL_DEVICE_ONLY
2249+
template <typename T, int N> struct VecStorageImpl {
2250+
using DataType = std::array<T, (N == 3) ? 4 : N>;
2251+
};
2252+
#endif // __SYCL_DEVICE_ONLY__
22692253
#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)
22702254

22712255
#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
@@ -2283,7 +2267,6 @@ template <typename T, int N> struct VecStorageImpl;
22832267
using DataType = ::cl_##cl_type##num; \
22842268
};
22852269
#endif // __SYCL_USE_EXT_VECTOR_TYPE__
2286-
#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
22872270

22882271
#ifndef __SYCL_USE_EXT_VECTOR_TYPE__
22892272
#define __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(type, cl_type) \
@@ -2307,12 +2290,15 @@ __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(double, double)
23072290
#undef __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE
23082291
#undef __SYCL_DEFINE_VECSTORAGE_IMPL
23092292
#endif // ifndef __SYCL_USE_EXT_VECTOR_TYPE__
2293+
#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
23102294

23112295
// Single element bool
23122296
template <> struct VecStorage<bool, 1, void> {
23132297
using DataType = bool;
23142298
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2299+
#ifdef __SYCL_DEVICE_ONLY__
23152300
using VectorDataType = bool;
2301+
#endif // __SYCL_DEVICE_ONLY__
23162302
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23172303
};
23182304

@@ -2324,41 +2310,65 @@ struct VecStorage<bool, N, typename std::enable_if_t<isValidVectorSize(N)>> {
23242310
std::int32_t, std::int64_t>,
23252311
N>::DataType;
23262312
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2313+
#ifdef __SYCL_DEVICE_ONLY__
23272314
using VectorDataType =
23282315
typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
23292316
std::int32_t, std::int64_t>,
23302317
N>::VectorDataType;
2318+
#endif // __SYCL_DEVICE_ONLY__
2319+
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
2320+
};
2321+
2322+
#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
2323+
// Single element byte. Multiple elements will propagate through a later
2324+
// specialization.
2325+
template <> struct VecStorage<std::byte, 1, void> {
2326+
using DataType = std::int8_t;
2327+
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2328+
#ifdef __SYCL_DEVICE_ONLY__
2329+
using VectorDataType = std::int8_t;
2330+
#endif // __SYCL_DEVICE_ONLY__
23312331
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23322332
};
2333+
#endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
23332334

23342335
// Single element signed integers
23352336
template <typename T>
23362337
struct VecStorage<T, 1, typename std::enable_if_t<is_sigeninteger_v<T>>> {
2338+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
23372339
using DataType = select_apply_cl_t<T, std::int8_t, std::int16_t, std::int32_t,
23382340
std::int64_t>;
2339-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2341+
#else // __INTEL_PREVIEW_BREAKING_CHANGES
2342+
using DataType = T;
2343+
#ifdef __SYCL_DEVICE_ONLY__
23402344
using VectorDataType = DataType;
2345+
#endif // __SYCL_DEVICE_ONLY__
23412346
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23422347
};
23432348

23442349
// Single element unsigned integers
23452350
template <typename T>
23462351
struct VecStorage<T, 1, typename std::enable_if_t<is_sugeninteger_v<T>>> {
2352+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
23472353
using DataType = select_apply_cl_t<T, std::uint8_t, std::uint16_t,
23482354
std::uint32_t, std::uint64_t>;
2349-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2355+
#else // __INTEL_PREVIEW_BREAKING_CHANGES
2356+
using DataType = T;
2357+
#ifdef __SYCL_DEVICE_ONLY__
23502358
using VectorDataType = DataType;
2359+
#endif // __SYCL_DEVICE_ONLY__
23512360
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23522361
};
23532362

23542363
// Single element floating-point (except half)
23552364
template <typename T>
23562365
struct VecStorage<
23572366
T, 1, typename std::enable_if_t<!is_half_v<T> && is_sgenfloat_v<T>>> {
2358-
using DataType =
2359-
select_apply_cl_t<T, std::false_type, std::false_type, float, double>;
2367+
using DataType = T;
23602368
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2369+
#ifdef __SYCL_DEVICE_ONLY__
23612370
using VectorDataType = DataType;
2371+
#endif // __SYCL_DEVICE_ONLY__
23622372
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23632373
};
23642374
// Multiple elements signed/unsigned integers and floating-point (except half)
@@ -2371,32 +2381,39 @@ struct VecStorage<
23712381
using DataType =
23722382
typename VecStorageImpl<typename VecStorage<T, 1>::DataType, N>::DataType;
23732383
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2384+
#ifdef __SYCL_DEVICE_ONLY__
23742385
using VectorDataType =
23752386
typename VecStorageImpl<typename VecStorage<T, 1>::DataType,
23762387
N>::VectorDataType;
2388+
#endif // __SYCL_DEVICE_ONLY__
23772389
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23782390
};
23792391

23802392
// Single element half
23812393
template <> struct VecStorage<half, 1, void> {
23822394
using DataType = sycl::detail::half_impl::StorageT;
23832395
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2396+
#ifdef __SYCL_DEVICE_ONLY__
23842397
using VectorDataType = sycl::detail::half_impl::StorageT;
2398+
#endif // __SYCL_DEVICE_ONLY__
23852399
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
23862400
};
2401+
23872402
// Multiple elements half
2388-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
2403+
#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) && defined(__SYCL_DEVICE_ONLY__)
23892404
#define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
23902405
template <> struct VecStorage<half, Num, void> { \
23912406
using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
23922407
using VectorDataType = sycl::detail::half_impl::Vec##Num##StorageT; \
23932408
};
2394-
#else // __INTEL_PREVIEW_BREAKING_CHANGES
2409+
#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) &&
2410+
// defined(__SYCL_DEVICE_ONLY__)
23952411
#define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
23962412
template <> struct VecStorage<half, Num, void> { \
23972413
using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
23982414
};
2399-
#endif
2415+
#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) &&
2416+
// defined(__SYCL_DEVICE_ONLY__)
24002417

24012418
__SYCL_DEFINE_HALF_VECSTORAGE(2)
24022419
__SYCL_DEFINE_HALF_VECSTORAGE(3)

0 commit comments

Comments
 (0)