[SYCL] vec abi unification and trivially copyable

sycl/include/sycl/detail/vector_traits.hpp

@@ -17,12 +17,20 @@ inline namespace _V1 {
 namespace detail {
 
 // 4.10.2.6 Memory layout and alignment
-template <typename T, int N>
+// due to MSVC the maximum alignment for sycl::vec is 64 and this proposed
+// change is being brought to the spec committee.
+constexpr size_t MaxVecAlignment = 64;
+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>> {};
+    : std::conditional_t<
+          N == 3,
+          std::integral_constant<size_t,
+                                 (std::min)(sizeof(T) * 4, MaxVecAlignment)>,
+          std::integral_constant<size_t,
+                                 (std::min)(sizeof(T) * N, MaxVecAlignment)>> {
+};
 
-template <typename T, int N>
+template <typename T, size_t N>
 struct vector_alignment
     : vector_alignment_impl<std::remove_cv_t<std::remove_reference_t<T>>, N> {};
 } // namespace detail

sycl/include/sycl/half_type.hpp

@@ -252,28 +252,12 @@ using BIsRepresentationT = half;
 // for vec because they are actually defined as an integer type under the
 // hood. As a result half values will be converted to the integer and passed
 // as a kernel argument which is expected to be floating point number.
-template <int NumElements> struct half_vec {
-  alignas(
-      vector_alignment<StorageT, NumElements>::value) StorageT s[NumElements];
 
-  __SYCL_CONSTEXPR_HALF half_vec() : s{0.0f} { initialize_data(); }
-  template <typename... Ts,
-            typename = std::enable_if_t<(sizeof...(Ts) == NumElements) &&
-                                        (std::is_same_v<half, Ts> && ...)>>
-  __SYCL_CONSTEXPR_HALF half_vec(const Ts &...hs) : s{hs...} {}
-
-  constexpr void initialize_data() {
-    for (size_t i = 0; i < NumElements; ++i) {
-      s[i] = StorageT(0.0f);
-    }
-  }
-};
-
-using Vec2StorageT = half_vec<2>;
-using Vec3StorageT = half_vec<3>;
-using Vec4StorageT = half_vec<4>;
-using Vec8StorageT = half_vec<8>;
-using Vec16StorageT = half_vec<16>;
+using Vec2StorageT = std::array<StorageT, 2>;
+using Vec3StorageT = std::array<StorageT, 3>;
+using Vec4StorageT = std::array<StorageT, 4>;
+using Vec8StorageT = std::array<StorageT, 8>;
+using Vec16StorageT = std::array<StorageT, 16>;
 #endif
 
 #ifndef __SYCL_DEVICE_ONLY__

sycl/source/CMakeLists.txt

@@ -92,6 +92,12 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME)
   else()
     target_compile_options(${LIB_OBJ_NAME} PUBLIC
                            -fvisibility=hidden -fvisibility-inlines-hidden)
+
+    # Sycl math built-in macros cause a GCC 4.6 'note' to be output repeatedly.
+    # => note: the ABI for passing parameters with 32-byte alignment has changed in GCC 4.6
+    # Seems to be no way to suppress it except use -Wno-psabi
+    target_compile_options(${LIB_OBJ_NAME} PUBLIC -Wno-psabi)
+
     if (UNIX AND NOT APPLE)
       set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/ld-version-script.txt")
       target_link_libraries(

sycl/source/detail/builtins_relational.cpp

@@ -458,9 +458,15 @@ __SYCL_EXPORT rel_res_t sycl_host_SignBitSet(s::cl_float x) __NOEXC {
 __SYCL_EXPORT rel_res_t sycl_host_SignBitSet(s::cl_double x) __NOEXC {
   return std::signbit(x);
 }
+#ifndef __GNUC__
+// GCC 11.3.0 (and friends) has an internal compiler error when this function is
+// passed to the macro declaration below. ( MAKE_1V ... 4 ... ) Fortunately, we
+// don't actually need it when compiling the SYCL library itself.
+// TODO: switch to templates instead of these overload expansions.
 __SYCL_EXPORT s::cl_int __vSignBitSet(s::cl_float x) __NOEXC {
   return -static_cast<s::cl_int>(std::signbit(x));
 }
+#endif
 __SYCL_EXPORT s::cl_long __vSignBitSet(s::cl_double x) __NOEXC {
   return -static_cast<s::cl_long>(std::signbit(x));
 }

sycl/test/abi/sycl_symbols_linux.dump

@@ -19,7 +19,6 @@ _ZN10__host_std11__vIsNormalEd
 _ZN10__host_std11__vIsNormalEf
 _ZN10__host_std13__vSignBitSetEN4sycl3_V16detail9half_impl4halfE
 _ZN10__host_std13__vSignBitSetEd
-_ZN10__host_std13__vSignBitSetEf
 _ZN10__host_std13sycl_host_AllEN4sycl3_V13vecIaLi16EEE
 _ZN10__host_std13sycl_host_AllEN4sycl3_V13vecIaLi1EEE
 _ZN10__host_std13sycl_host_AllEN4sycl3_V13vecIaLi2EEE

sycl/test/basic_tests/types.cpp

@@ -21,14 +21,17 @@ template <typename T, int N> inline void checkVectorSizeAndAlignment() {
   using VectorT = s::vec<T, N>;
   constexpr auto RealLength = (N != 3 ? N : 4);
   static_assert(sizeof(VectorT) == (sizeof(T) * RealLength), "");
-#if defined(_WIN32) && (_MSC_VER) &&                                           \
-    defined(__NO_EXT_VECTOR_TYPE_ON_HOST__) && !defined(__SYCL_DEVICE_ONLY__)
-  // See comments around __SYCL_ALIGNED_VAR macro definition in types.hpp
-  // We can't enforce proper alignment of "huge" vectors (>64 bytes) on Windows
-  // and the test exposes this limitation.
-  if constexpr (alignof(T) * RealLength < 64)
-#endif
+
+  // SYCL 2020 spec says that alignment is supposed to be same as size,
+  // but MSVC won't allow an alignment of anything larger than 64 for
+  // a direct parameter. The math built-ins use direct param calls.
+  // It has been decided to change the spec to have a max alignment of
+  // 64.
+  if constexpr (alignof(T) * RealLength <= 64)
     static_assert(alignof(VectorT) == (alignof(T) * RealLength), "");
+  else
+    static_assert(alignof(VectorT) == 64,
+                  "huge vectors should have a maximum alignment of 64");
 }
 
 template <typename T> inline void checkVectorsWithN() {

sycl/test/basic_tests/valid_kernel_args.cpp

@@ -35,10 +35,11 @@ template <typename T> void check() {
 }
 
 SYCL_EXTERNAL void foo() {
-#ifdef __SYCL_DEVICE_ONLY__
+
   check<int>();
   check<sycl::vec<sycl::opencl::cl_uchar, 4>>();
   check<SomeStructure>();
-#endif
+  check<sycl::int4>();
+  check<sycl::long16>();
   check<SomeMarrayStructure>();
 }