diff --git a/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp b/sycl/include/sycl/detail/type_traits/vec_marray_traits.hpp new file mode 100644 index 0000000000000..6ce39bf6a072a --- /dev/null +++ b/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 +#include + +#include + +namespace sycl { +inline namespace _V1 { +template class __SYCL_EBO vec; + +template class marray; + +namespace detail { +template 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 struct simplify_if_swizzle { + using type = T; +}; + +template class OperationCurrentT, int... Indexes> +struct simplify_if_swizzle> { + using type = vec; +}; + +template +using simplify_if_swizzle_t = typename simplify_if_swizzle::type; + +// --------- is_* traits ------------------ // +template struct is_vec : std::false_type {}; +template struct is_vec> : std::true_type {}; +template constexpr bool is_vec_v = is_vec::value; + +template +struct is_ext_vector : std::false_type {}; +template +struct is_valid_type_for_ext_vector : std::false_type {}; +#if defined(__has_extension) +#if __has_extension(attribute_ext_vector_type) +template +using ext_vector = T __attribute__((ext_vector_type(N))); +template +struct is_ext_vector> : std::true_type {}; +template +struct is_valid_type_for_ext_vector>> + : std::true_type {}; +#endif +#endif +template +inline constexpr bool is_ext_vector_v = is_ext_vector::value; +template +inline constexpr bool is_valid_type_for_ext_vector_v = + is_valid_type_for_ext_vector::value; + +template struct is_swizzle : std::false_type {}; +template class OperationCurrentT, int... Indexes> +struct is_swizzle> : std::true_type {}; +template constexpr bool is_swizzle_v = is_swizzle::value; + +template +constexpr bool is_vec_or_swizzle_v = is_vec_v || is_swizzle_v; + +template struct is_marray : std::false_type {}; +template +struct is_marray> : std::true_type {}; +template constexpr bool is_marray_v = is_marray::value; + +// --------- num_elements trait ------------------ // +template +struct num_elements : std::integral_constant {}; +template +struct num_elements> : std::integral_constant {}; +template +struct num_elements> + : std::integral_constant {}; +#if defined(__has_extension) +#if __has_extension(attribute_ext_vector_type) +template +struct num_elements + : std::integral_constant {}; +#endif +#endif +template class OperationCurrentT, int... Indexes> +struct num_elements> + : std::integral_constant {}; + +template +inline constexpr std::size_t num_elements_v = num_elements::value; + +// --------- element_type trait ------------------ // +template struct element_type { + using type = T; +}; +template struct element_type> { + using type = T; +}; +template struct element_type> { + using type = T; +}; +#if defined(__has_extension) +#if __has_extension(attribute_ext_vector_type) +template +struct element_type { + using type = T; +}; +#endif +#endif +template using element_type_t = typename element_type::type; + +template +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 diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp index ea935032ba445..1d7e2cdd44926 100644 --- a/sycl/include/sycl/vector.hpp +++ b/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 class vec_base_test; + template 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 should be equal to float16 in memory. - using DataType = std::array; + // + // 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) == sizeof(DataT[AdjustedNum]) && + alignof(std::array) == + alignof(DataT[AdjustedNum]), +#endif + DataT[AdjustedNum], std::array>; + + // To allow testing of private methods + template friend class detail::vec_base_test; #ifdef __SYCL_DEVICE_ONLY__ using element_type_for_vector_t = typename detail::map_type< diff --git a/sycl/test/abi/layout_vec.cpp b/sycl/test/abi/layout_vec.cpp index 1f61d0fcd4666..06a9a7959a530 100644 --- a/sycl/test/abi/layout_vec.cpp +++ b/sycl/test/abi/layout_vec.cpp @@ -12,8 +12,7 @@ SYCL_EXTERNAL void foo(sycl::vec) {} // CHECK: 0 | class sycl::vec // ignore empty base classes -// CHECK: 0 | struct std::array 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) {} // 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] diff --git a/sycl/test/basic_tests/vectors/storage.cpp b/sycl/test/basic_tests/vectors/storage.cpp new file mode 100644 index 0000000000000..bbb14cfe24227 --- /dev/null +++ b/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 + +#include + +namespace sycl { +namespace detail { +template class vec_base_test { +public: + static void do_check() { + constexpr bool uses_std_array = + std::is_same_v::DataType, std::array>; + constexpr bool uses_plain_array = + std::is_same_v::DataType, T[N]>; + + constexpr bool std_array_and_plain_array_have_the_same_layout = + sizeof(std::array) == sizeof(T[N]) && + alignof(std::array) == 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::do_check(); } diff --git a/sycl/test/regression/vec_array_windows.cpp b/sycl/test/regression/vec_array_windows.cpp new file mode 100644 index 0000000000000..d9511de372f54 --- /dev/null +++ b/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 + +// 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 v) { return v[0]; }