Skip to content

[NFC][SYCL] vec code unification #11879

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 3 commits into from
Closed

[NFC][SYCL] vec code unification #11879

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
31d778d
[SYCL][NFC] Unify vec implementation
AlexeySachkov Nov 14, 2023
100e1e0
[SYCL][NFC] More unification
AlexeySachkov Nov 14, 2023
7ae01c4
Try to fix alignment setting
AlexeySachkov Nov 14, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
38 changes: 30 additions & 8 deletions sycl/include/sycl/detail/vector_traits.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@
#pragma once

#include <algorithm> // for std::min and vs2017 win
#include <limits> // for numeric_limits
#include <type_traits> // for integral_constant, conditional_t, remove_cv_t

namespace sycl {
Expand All @@ -17,9 +18,36 @@ namespace detail {

// 4.10.2.6 Memory layout and alignment
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
// due to MSVC the maximum alignment for sycl::vec is 64 and this proposed
// change is being brought to the spec committee.
// SYCL 2020 vec alignment requirements have been relaxed in
// KhronosGroup/SYCL-Docs#448. New specification wording only guarantees 64-byte
// alignment of vec class and we leverage this here to avoid dealing with MSVC
// limitations (see below).
constexpr size_t MaxVecAlignment = 64;
#else
// This version is preserved to maintain API/ABI compatibility with older
// releases.
// FIXME: drop this branch once API/ABI break is allowed

#if defined(_WIN32) && (_MSC_VER)
// MSVC Compiler doesn't allow using of function arguments with alignment
// requirements. MSVC Compiler Error C2719: 'parameter': formal parameter with
// __declspec(align('#')) won't be aligned. The align __declspec modifier
// is not permitted on function parameters. Function parameter alignment
// is controlled by the calling convention used.
// For more information, see Calling Conventions
// (https://docs.microsoft.com/en-us/cpp/cpp/calling-conventions).
// For information on calling conventions for x64 processors, see
// Calling Convention
// (https://docs.microsoft.com/en-us/cpp/build/x64-calling-convention).
constexpr size_t MaxVecAlignment = 64;
#else
// To match ABI of previos releases, we don't impose any restrictions on vec
// alignment on Linux
constexpr size_t MaxVecAlignment = std::numeric_limits<size_t>::max();
#endif

#endif // __INTEL_PREVIEW_BREAKING_CHANGES

template <typename T, size_t N>
struct vector_alignment_impl
: std::conditional_t<
Expand All @@ -29,12 +57,6 @@ struct vector_alignment_impl
std::integral_constant<size_t,
(std::min)(sizeof(T) * N, MaxVecAlignment)>> {
};
#else
template <typename T, size_t N>
struct vector_alignment_impl
: std::conditional_t<N == 3, std::integral_constant<int, sizeof(T) * 4>,
std::integral_constant<int, sizeof(T) * N>> {};
#endif

template <typename T, size_t N>
struct vector_alignment
Expand Down
94 changes: 26 additions & 68 deletions sycl/include/sycl/types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -272,30 +272,6 @@ template <typename T> using vec_data = detail::vec_helper<T>;
template <typename T>
using vec_data_t = typename detail::vec_helper<T>::RetType;

#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)

#if defined(_WIN32) && (_MSC_VER)
// MSVC Compiler doesn't allow using of function arguments with alignment
// requirements. MSVC Compiler Error C2719: 'parameter': formal parameter with
// __declspec(align('#')) won't be aligned. The align __declspec modifier
// is not permitted on function parameters. Function parameter alignment
// is controlled by the calling convention used.
// For more information, see Calling Conventions
// (https://docs.microsoft.com/en-us/cpp/cpp/calling-conventions).
// For information on calling conventions for x64 processors, see
// Calling Convention
// (https://docs.microsoft.com/en-us/cpp/build/x64-calling-convention).
#pragma message("Alignment of class vec is not in accordance with SYCL \
specification requirements, a limitation of the MSVC compiler(Error C2719).\
Requested alignment applied, limited at 64.")
#define __SYCL_ALIGNED_VAR(type, x, var) \
type __declspec(align((x < 64) ? x : 64)) var
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This doesn't seem right. You are now returning us to having Windows use __declspec(align(...)) in both the -fpreview and not.
Part of the problem is that __declspec(align(...)) is not respected by the device at all. It means nothing. To get the alignments to agree, we need to use alignas on both host and device on both linux and windows when using the -fpreview flag. When not using the -fpreview flag, we leave things as they were.

#else
#define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
Comment on lines -291 to -294
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ok, I see now why we need that macro :) Please expect another attempt at unifying this in the next commit

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

7ae01c4

#endif

#endif //! defined(__INTEL_PREVIEW_BREAKING_CHANGES)

/// Provides a cross-patform vector class template that works efficiently on
/// SYCL devices as well as in host C++ code.
///
Expand Down Expand Up @@ -1458,30 +1434,33 @@ template <typename Type, int NumElements> class vec {
return (NumElements == 1) ? getValue(Index, 0) : getValue(Index, 0.f);
}

#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)

// fields

// Alignment is the same as size, to a maximum size of 64.
// detail::vector_alignment will return that value.
alignas(detail::vector_alignment<DataT, NumElements>::value) DataType m_Data;

#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES)

#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)

// fields

// Used "__SYCL_ALIGNED_VAR" instead "alignas" to handle MSVC compiler.
// For MSVC compiler max alignment is 64, e.g. vec<double, 16> required
// alignment of 128 and MSVC compiler cann't align a parameter with requested
// alignment of 128. For alignment request larger than 64, 64-alignment
// is applied
// Alignment is the same as size, to a maximum size of 64 (with some
// exceptions, see detail::vector_alignment).
#if defined(_WIN32) && (_MSC_VER)
#define __SYCL_ALIGNED_VAR(type, x, var) \
type __declspec(align((x < MaxVecAlignment) ? x : MaxVecAlignment)) var
#else
#define __SYCL_ALIGNED_VAR(type, x, var) alignas(x) type var
#endif
// Used "__SYCL_ALIGNED_VAR" instead of "alignas" to handle MSVC compiler.
//
// SYCL 2020 spec allows us to have at most 64-byte alignment for vec, but
// alignas requires that passed alignment is greater or equal to a minumum
// required alignment for a type. That is 128 bytes for types like
// vec<double, 16> and MSVC is not able to support that, which makes it
// impossible to use alignas directly.
//
// We have prepared a vec refactoring which changes underlying storage data
// type so we are able to use alignas directly, but it is hidden under preview
// breaking changes macro for now.
// FIXME: we should be able to drop the macro and directly use alignas once
// functionality under preview breaking changes macro is promoted.
__SYCL_ALIGNED_VAR(DataType,
(detail::vector_alignment<DataT, NumElements>::value),
m_Data);

#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
#undef __SYCL_ALIGNED_VAR

// friends
template <typename T1, typename T2, typename T3, template <typename> class T4,
Expand Down Expand Up @@ -2273,6 +2252,8 @@ template <typename T, int N> struct VecStorageImpl;
#ifdef __SYCL_USE_EXT_VECTOR_TYPE__
template <typename T, int N> struct VecStorageImpl {
using DataType = T __attribute__((ext_vector_type(N)));
using VectorDataType =
DataType; // to unify code with preview breaking changes mode
};
#else
// When ext_vector_type is not available, we rely on cl_* types from CL/cl.h
Expand All @@ -2281,6 +2262,8 @@ template <typename T, int N> struct VecStorageImpl;
#define __SYCL_DEFINE_VECSTORAGE_IMPL(type, cl_type, num) \
template <> struct VecStorageImpl<type, num> { \
using DataType = ::cl_##cl_type##num; \
using VectorDataType = \
DataType; /* to unify code with preview breaking changes mode */ \
};
#endif // __SYCL_USE_EXT_VECTOR_TYPE__
#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
Expand Down Expand Up @@ -2311,9 +2294,7 @@ __SYCL_DEFINE_VECSTORAGE_IMPL_FOR_TYPE(double, double)
// Single element bool
template <> struct VecStorage<bool, 1, void> {
using DataType = bool;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType = bool;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};

// Multiple element bool
Expand All @@ -2323,32 +2304,26 @@ struct VecStorage<bool, N, typename std::enable_if_t<isValidVectorSize(N)>> {
typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
std::int32_t, std::int64_t>,
N>::DataType;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType =
typename VecStorageImpl<select_apply_cl_t<bool, std::int8_t, std::int16_t,
std::int32_t, std::int64_t>,
N>::VectorDataType;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};

// Single element signed integers
template <typename T>
struct VecStorage<T, 1, typename std::enable_if_t<is_sigeninteger_v<T>>> {
using DataType = select_apply_cl_t<T, std::int8_t, std::int16_t, std::int32_t,
std::int64_t>;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType = DataType;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};

// Single element unsigned integers
template <typename T>
struct VecStorage<T, 1, typename std::enable_if_t<is_sugeninteger_v<T>>> {
using DataType = select_apply_cl_t<T, std::uint8_t, std::uint16_t,
std::uint32_t, std::uint64_t>;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType = DataType;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};

// Single element floating-point (except half)
Expand All @@ -2357,9 +2332,7 @@ struct VecStorage<
T, 1, typename std::enable_if_t<!is_half_v<T> && is_sgenfloat_v<T>>> {
using DataType =
select_apply_cl_t<T, std::false_type, std::false_type, float, double>;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType = DataType;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};
// Multiple elements signed/unsigned integers and floating-point (except half)
template <typename T, int N>
Expand All @@ -2370,33 +2343,22 @@ struct VecStorage<
(is_sgenfloat_v<T> && !is_half_v<T>))>> {
using DataType =
typename VecStorageImpl<typename VecStorage<T, 1>::DataType, N>::DataType;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType =
typename VecStorageImpl<typename VecStorage<T, 1>::DataType,
N>::VectorDataType;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};

// Single element half
template <> struct VecStorage<half, 1, void> {
using DataType = sycl::detail::half_impl::StorageT;
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
using VectorDataType = sycl::detail::half_impl::StorageT;
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
};
// Multiple elements half
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
#define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
template <> struct VecStorage<half, Num, void> { \
using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
using VectorDataType = sycl::detail::half_impl::Vec##Num##StorageT; \
};
#else // __INTEL_PREVIEW_BREAKING_CHANGES
#define __SYCL_DEFINE_HALF_VECSTORAGE(Num) \
template <> struct VecStorage<half, Num, void> { \
using DataType = sycl::detail::half_impl::Vec##Num##StorageT; \
};
#endif

__SYCL_DEFINE_HALF_VECSTORAGE(2)
__SYCL_DEFINE_HALF_VECSTORAGE(3)
Expand Down Expand Up @@ -2554,7 +2516,3 @@ struct CheckDeviceCopyable<

} // namespace _V1
} // namespace sycl

#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES)
#undef __SYCL_ALIGNED_VAR
#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES)