Skip to content

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

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

Open
wants to merge 2 commits into
base: sycl
Choose a base branch
from
Open

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4e1ad13
[SYCL] Work around codegen issue with vectors-of-3
hvdijk Jun 17, 2025
e5c2a0d
Use bit_cast for other sizes, fix another instance.
hvdijk Jun 25, 2025
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
17 changes: 15 additions & 2 deletions sycl/include/sycl/detail/builtins/helper_macros.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)); \
} \
}

Expand Down
16 changes: 13 additions & 3 deletions sycl/include/sycl/detail/builtins/math_functions.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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__
Expand Down
5 changes: 4 additions & 1 deletion sycl/include/sycl/vector.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -527,12 +527,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
Expand Down