[SYCL] Work around codegen issue with vectors-of-3 (#19087)

Vectors of length 3 have an adjusted length of 4 where the last element
is poison. Copying that poison element results in wrong code being
generated. Avoid that by only copying the first 3 elements.

Author: hvdijk

sycl/include/sycl/detail/builtins/helper_macros.hpp

@@ -192,13 +192,26 @@
   template <NUM_ARGS##_TYPENAME_TYPE>                                          \
   detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>(NAME)(                             \
       NUM_ARGS##_TEMPLATE_TYPE_ARG) {                                          \
+    using ToTy = detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>;                    \
     if constexpr (detail::is_marray_v<T0>) {                                   \
       return detail::DELEGATOR(                                                \
           [](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); },      \
           NUM_ARGS##_ARG);                                                     \
+    } else if constexpr (detail::is_vec_v<ToTy>) {                             \
+      if constexpr (ToTy::size() == 3) {                                       \
+        /* For vectors of length 3, make sure to only copy 3 elements, not 4,  \
+           to work around code generation issues, see LLVM #144454. */         \
+        auto From = __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG);                     \
+        ToTy To;                                                               \
+        constexpr auto N =                                                     \
+            ToTy::size() * sizeof(detail::get_elem_type_t<ToTy>);              \
+        sycl::detail::memcpy_no_adl(&To, &From, N);                            \
+        return To;                                                             \
+      } else {                                                                 \
+        return bit_cast<ToTy>(__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));          \
+      }                                                                        \
     } else {                                                                   \
-      return bit_cast<detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>>(              \
-          __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));                              \
+      return bit_cast<ToTy>(__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG));            \
     }                                                                          \
   }
 

sycl/include/sycl/detail/builtins/math_functions.inc

@@ -254,9 +254,19 @@ auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) {
   detail::NON_SCALAR_ENABLER<SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE),  \
                              PtrTy>                                            \
   NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), PtrTy p) {         \
-    return bit_cast<detail::NON_SCALAR_ENABLER<                                \
-        SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), PtrTy>>(              \
-        detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p));        \
+    auto From = detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
+    using ToTy = detail::NON_SCALAR_ENABLER<                                   \
+        SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), PtrTy>;               \
+    if constexpr (ToTy::size() == 3) {                                         \
+      /* For vectors of length 3, make sure to only copy 3 elements, not 4, to \
+         work around code generation issues, see LLVM #144454. */              \
+      ToTy To;                                                                 \
+      constexpr auto N = ToTy::size() * sizeof(detail::get_elem_type_t<ToTy>); \
+      sycl::detail::memcpy_no_adl(&To, &From, N);                              \
+      return To;                                                               \
+    } else {                                                                   \
+      return bit_cast<ToTy>(From);                                             \
+    }                                                                          \
   }
 
 #if __SYCL_DEVICE_ONLY__

sycl/include/sycl/vector.hpp

@@ -530,12 +530,15 @@ class __SYCL_EBO vec :
   // `vector_t` is the same as `DataT`. Not that the other ctor isn't a template
   // so we don't even need a smart `enable_if` condition here, the mere fact of
   // this being a template makes the other ctor preferred.
+  // For vectors of length 3, make sure to only copy 3 elements, not 4, to work
+  // around code generation issues, see LLVM #144454.
   template <
       typename vector_t_ = vector_t,
       typename = typename std::enable_if_t<std::is_same_v<vector_t_, vector_t>>>
   constexpr vec(vector_t_ openclVector) {
     sycl::detail::memcpy_no_adl(&this->m_Data, &openclVector,
-                                sizeof(openclVector));
+                                NumElements *
+                                    sizeof(element_type_for_vector_t));
   }
 
   /* @SYCL2020