Skip to content

[SYCL] Compile time/remove vec convert from builtins #21698

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

Merged
KornevNikita merged 4 commits into from
Apr 16, 2026
Merged

[SYCL] Compile time/remove vec convert from builtins #21698

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
a8004be
[SYCL] Decouple builtins from vec::convert
koparasy Apr 5, 2026
b6367fd
Fix formatting
koparasy Apr 10, 2026
66bfcb5
Address comments
koparasy Apr 14, 2026
7223b14
Address comments
koparasy Apr 14, 2026
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
29 changes: 28 additions & 1 deletion sycl/include/sycl/detail/builtins/builtins.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -64,10 +64,11 @@
#pragma once

#include <sycl/detail/fwd/multi_ptr.hpp>
#include <sycl/detail/generic_type_traits.hpp>
#include <sycl/detail/helpers.hpp>
#include <sycl/detail/type_traits.hpp>
#include <sycl/detail/type_traits/vec_marray_traits.hpp>
#include <sycl/detail/vector_convert.hpp>
#include <sycl/half_type.hpp>
#include <sycl/marray.hpp>
#include <sycl/vector.hpp>

Expand Down Expand Up @@ -203,6 +204,32 @@ template <class T, int N> marray<T, N> to_marray(vec<T, N> X) {
return Marray;
}

// Relation builtins widen signed-char masks to the required integer element
// type. Keep that conversion local here so builtins.hpp does not need to pull
// in vector_convert.hpp just for vec::convert.
template <typename NewElemT, int N>
vec<NewElemT, N> relational_mask_widen(vec<signed char, N> X) {
static_assert(std::is_integral_v<NewElemT> &&
!std::is_same_v<NewElemT, bool>);

#if defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__)
// Keep NVPTX on the scalar fallback for consistency with vec::convert.
// TODO: Likely unnecessary as https://github.com/intel/llvm/issues/11840
// has been closed already.
if constexpr (N > 1) {
using src_vector_t = signed char __attribute__((ext_vector_type(N)));
using dst_vector_t = NewElemT __attribute__((ext_vector_type(N)));
auto OpenCLVec = bit_cast<src_vector_t>(X);
return bit_cast<vec<NewElemT, N>>(
__builtin_convertvector(OpenCLVec, dst_vector_t));
}
#endif // defined(__SYCL_DEVICE_ONLY__) && !defined(__NVPTX__)

vec<NewElemT, N> Result{};
loop<N>([&](auto idx) { Result[idx] = static_cast<NewElemT>(X[idx]); });
return Result;
}
Comment thread
uditagarwal97 marked this conversation as resolved.

namespace builtins {
#ifdef __SYCL_DEVICE_ONLY__
template <typename T> auto convert_arg(T &&x) {
Expand Down
2 changes: 1 addition & 1 deletion sycl/include/sycl/detail/builtins/relational_functions.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,7 +66,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) {
auto tmp = bit_cast<vec<signed char, num_elements<T>::value>>(ret);
using res_elem_type = fixed_width_signed<sizeof(get_elem_type_t<T>)>;
static_assert(is_scalar_arithmetic_v<res_elem_type>);
return tmp.template convert<res_elem_type>();
return relational_mask_widen<res_elem_type>(tmp);
} else if constexpr (std::is_same_v<T, half>) {
return bool{F(builtins::convert_arg(xs)...)};
} else {
Expand Down
10 changes: 5 additions & 5 deletions sycl/include/sycl/detail/image_accessor_util.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,12 +18,12 @@
#include <sycl/detail/array.hpp> // for array
#include <sycl/detail/export.hpp> // for __SYCL_EXPORT
#include <sycl/detail/generic_type_traits.hpp> // for max_v, min_v, TryToGe...
#include <sycl/detail/vector_convert.hpp> // for vec, vec::convert, operator*, round...
#include <sycl/exception.hpp>
#include <sycl/id.hpp> // for id
#include <sycl/image.hpp> // for image_channel_type
#include <sycl/range.hpp> // for range
#include <sycl/sampler.hpp> // for addressing_mode, coor...
#include <sycl/vector.hpp> // for vec, operator*, round...
#include <sycl/id.hpp> // for id
#include <sycl/image.hpp> // for image_channel_type
#include <sycl/range.hpp> // for range
#include <sycl/sampler.hpp> // for addressing_mode, coor...

#include <cstdint> // for int32_t, uint16_t
#include <stddef.h> // for size_t
Expand Down
1 change: 1 addition & 0 deletions sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
#include <sycl/detail/builtins/builtins.hpp>
#include <sycl/detail/generic_type_traits.hpp>
#include <sycl/detail/type_traits.hpp>
#include <sycl/exception.hpp>

namespace sycl {
inline namespace _V1 {
Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/Basic/char_builtins.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
#include <array>
#include <sycl/builtins.hpp>
#include <sycl/detail/core.hpp>
#include <sycl/detail/vector_convert.hpp>

using namespace sycl;

Expand Down
1 change: 1 addition & 0 deletions sycl/test-e2e/Basic/half_builtins.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
#include <sycl/detail/core.hpp>

#include <sycl/builtins.hpp>
#include <sycl/detail/vector_convert.hpp>

#include <cmath>
#include <limits>
Expand Down
14 changes: 7 additions & 7 deletions sycl/test/check_device_code/vector/bf16_builtins_new_vec.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -69,7 +69,7 @@ SYCL_EXTERNAL auto TestFMin(vec<bfloat16, 2> a, vec<bfloat16, 2> b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z8TestFMaxN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEES5_(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.70") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.70") align 8 captures(none) [[A:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.70") align 8 captures(none) [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 8 captures(none) [[A:%.*]], ptr noundef readonly byval({{.*}}) align 8 captures(none) [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I14_I:%.*]] = alloca <3 x float>, align 16
// CHECK-NEXT: [[DST_I_I_I_I15_I:%.*]] = alloca [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down Expand Up @@ -126,7 +126,7 @@ SYCL_EXTERNAL auto TestFMax(vec<bfloat16, 3> a, vec<bfloat16, 3> b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z9TestIsNanN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi4EEE(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.146") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.183") align 8 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 8 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I_I:%.*]] = alloca <4 x i16>, align 8
// CHECK-NEXT: [[DST_I_I_I_I_I:%.*]] = alloca [4 x float], align 4
Expand All @@ -142,8 +142,8 @@ SYCL_EXTERNAL auto TestFMax(vec<bfloat16, 3> a, vec<bfloat16, 3> b) {
// CHECK-NEXT: call void @llvm.lifetime.end.p0(ptr nonnull [[VEC_ADDR_I_I_I_I_I]]), !noalias [[META41]]
// CHECK-NEXT: call void @llvm.lifetime.end.p0(ptr nonnull [[DST_I_I_I_I_I]]), !noalias [[META41]]
// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = call spir_func noundef <4 x i8> @_Z13__spirv_IsNanDv4_f(<4 x float> noundef [[TMP1]]) #[[ATTR6]]
// CHECK-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = call spir_func noundef <4 x i32> @_Z22__spirv_SConvert_Rint4Dv4_a(<4 x i8> noundef [[CALL_I_I_I_I]]) #[[ATTR6]]
// CHECK-NEXT: [[CALL_I_I_I2_I:%.*]] = call spir_func noundef <4 x i16> @_Z24__spirv_SConvert_Rshort4Dv4_i(<4 x i32> noundef [[CALL_I_I_I_I_I_I]]) #[[ATTR6]]
// CHECK-NEXT: [[MASK_I32:%.*]] = sext <4 x i8> [[CALL_I_I_I_I]] to <4 x i32>
// CHECK-NEXT: [[CALL_I_I_I2_I:%.*]] = call spir_func noundef <4 x i16> @_Z24__spirv_SConvert_Rshort4Dv4_i(<4 x i32> noundef [[MASK_I32]]) #[[ATTR6]]
// CHECK-NEXT: store <4 x i16> [[CALL_I_I_I2_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !alias.scope [[META44:![0-9]+]]
// CHECK-NEXT: ret void
//
Expand All @@ -152,7 +152,7 @@ SYCL_EXTERNAL auto TestIsNan(vec<bfloat16, 4> a) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z8TestFabsN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.335") align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.335") align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32
// CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down Expand Up @@ -188,7 +188,7 @@ SYCL_EXTERNAL auto TestFabs(vec<bfloat16, 8> a) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z8TestCeilN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.335") align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.335") align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32
// CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down Expand Up @@ -224,7 +224,7 @@ SYCL_EXTERNAL auto TestCeil(vec<bfloat16, 8> a) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z7TestFMAN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi16EEES5_S5_(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.411") align 32 captures(none) initializes((0, 32)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.411") align 32 captures(none) [[A:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.411") align 32 captures(none) [[B:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.411") align 32 captures(none) [[C:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 32 captures(none) initializes((0, 32)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 32 captures(none) [[A:%.*]], ptr noundef readonly byval({{.*}}) align 32 captures(none) [[B:%.*]], ptr noundef readonly byval({{.*}}) align 32 captures(none) [[C:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I14_I:%.*]] = alloca <16 x float>, align 64
// CHECK-NEXT: [[DST_I_I_I_I15_I:%.*]] = alloca [16 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down
14 changes: 7 additions & 7 deletions sycl/test/check_device_code/vector/bf16_builtins_old_vec.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -69,7 +69,7 @@ SYCL_EXTERNAL auto TestFMin(vec<bfloat16, 2> a, vec<bfloat16, 2> b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z8TestFMaxN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi3EEES5_(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.70") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.70") align 8 captures(none) [[A:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.70") align 8 captures(none) [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 8 captures(none) [[A:%.*]], ptr noundef readonly byval({{.*}}) align 8 captures(none) [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I14_I:%.*]] = alloca <3 x float>, align 16
// CHECK-NEXT: [[DST_I_I_I_I15_I:%.*]] = alloca [4 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down Expand Up @@ -126,7 +126,7 @@ SYCL_EXTERNAL auto TestFMax(vec<bfloat16, 3> a, vec<bfloat16, 3> b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z9TestIsNanN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi4EEE(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.146") align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.184") align 8 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 8 captures(none) initializes((0, 8)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 8 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I_I:%.*]] = alloca <4 x i16>, align 8
// CHECK-NEXT: [[DST_I_I_I_I_I:%.*]] = alloca [4 x float], align 4
Expand All @@ -142,8 +142,8 @@ SYCL_EXTERNAL auto TestFMax(vec<bfloat16, 3> a, vec<bfloat16, 3> b) {
// CHECK-NEXT: call void @llvm.lifetime.end.p0(ptr nonnull [[VEC_ADDR_I_I_I_I_I]]), !noalias [[META41]]
// CHECK-NEXT: call void @llvm.lifetime.end.p0(ptr nonnull [[DST_I_I_I_I_I]]), !noalias [[META41]]
// CHECK-NEXT: [[CALL_I_I_I_I:%.*]] = call spir_func noundef <4 x i8> @_Z13__spirv_IsNanDv4_f(<4 x float> noundef [[TMP1]]) #[[ATTR6]]
// CHECK-NEXT: [[CALL_I_I_I_I_I_I:%.*]] = call spir_func noundef <4 x i32> @_Z22__spirv_SConvert_Rint4Dv4_a(<4 x i8> noundef [[CALL_I_I_I_I]]) #[[ATTR6]]
// CHECK-NEXT: [[CALL_I_I_I2_I:%.*]] = call spir_func noundef <4 x i16> @_Z24__spirv_SConvert_Rshort4Dv4_i(<4 x i32> noundef [[CALL_I_I_I_I_I_I]]) #[[ATTR6]]
// CHECK-NEXT: [[MASK_I32:%.*]] = sext <4 x i8> [[CALL_I_I_I_I]] to <4 x i32>
// CHECK-NEXT: [[CALL_I_I_I2_I:%.*]] = call spir_func noundef <4 x i16> @_Z24__spirv_SConvert_Rshort4Dv4_i(<4 x i32> noundef [[MASK_I32]]) #[[ATTR6]]
// CHECK-NEXT: store <4 x i16> [[CALL_I_I_I2_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !alias.scope [[META44:![0-9]+]]
// CHECK-NEXT: ret void
//
Expand All @@ -152,7 +152,7 @@ SYCL_EXTERNAL auto TestIsNan(vec<bfloat16, 4> a) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z8TestFabsN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.336") align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.336") align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32
// CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down Expand Up @@ -188,7 +188,7 @@ SYCL_EXTERNAL auto TestFabs(vec<bfloat16, 8> a) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z8TestCeilN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi8EEE(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.336") align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.336") align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 16 captures(none) initializes((0, 16)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 16 captures(none) [[A:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I2_I:%.*]] = alloca <8 x float>, align 32
// CHECK-NEXT: [[DST_I_I_I_I3_I:%.*]] = alloca [8 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down Expand Up @@ -224,7 +224,7 @@ SYCL_EXTERNAL auto TestCeil(vec<bfloat16, 8> a) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z7TestFMAN4sycl3_V13vecINS0_3ext6oneapi8bfloat16ELi16EEES5_S5_(
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret(%"class.sycl::_V1::vec.412") align 32 captures(none) initializes((0, 32)) [[AGG_RESULT:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.412") align 32 captures(none) [[A:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.412") align 32 captures(none) [[B:%.*]], ptr noundef readonly byval(%"class.sycl::_V1::vec.412") align 32 captures(none) [[C:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias writable writeonly sret({{.*}}) align 32 captures(none) initializes((0, 32)) [[AGG_RESULT:%.*]], ptr noundef readonly byval({{.*}}) align 32 captures(none) [[A:%.*]], ptr noundef readonly byval({{.*}}) align 32 captures(none) [[B:%.*]], ptr noundef readonly byval({{.*}}) align 32 captures(none) [[C:%.*]]) local_unnamed_addr #[[ATTR0]] {{.*}}{
// CHECK-NEXT: entry:
// CHECK-NEXT: [[VEC_ADDR_I_I_I_I14_I:%.*]] = alloca <16 x float>, align 64
// CHECK-NEXT: [[DST_I_I_I_I15_I:%.*]] = alloca [16 x %"class.sycl::_V1::ext::oneapi::bfloat16"], align 2
Expand Down
20 changes: 9 additions & 11 deletions sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,31 +12,29 @@
// CHECK-NEXT: detail/builtins/builtins.hpp
// CHECK-NEXT: detail/fwd/multi_ptr.hpp
// CHECK-NEXT: access/access.hpp
// CHECK-NEXT: detail/generic_type_traits.hpp
// CHECK-NEXT: aliases.hpp
// CHECK-NEXT: bit_cast.hpp
// CHECK-NEXT: detail/fwd/half.hpp
// CHECK-NEXT: detail/type_traits.hpp
// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
// CHECK-NEXT: detail/helpers.hpp
// CHECK-NEXT: __spirv/spirv_types.hpp
// CHECK-NEXT: detail/defines.hpp
// CHECK-NEXT: detail/export.hpp
// CHECK-NEXT: memory_enums.hpp
// CHECK-NEXT: __spirv/spirv_vars.hpp
// CHECK-NEXT: detail/type_traits.hpp
// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
// CHECK-NEXT: detail/vector_convert.hpp
// CHECK-NEXT: detail/generic_type_traits.hpp
// CHECK-NEXT: aliases.hpp
// CHECK-NEXT: bit_cast.hpp
// CHECK-NEXT: detail/fwd/half.hpp
// CHECK-NEXT: detail/memcpy.hpp
// CHECK-NEXT: ext/oneapi/bfloat16.hpp
// CHECK-NEXT: half_type.hpp
// CHECK-NEXT: aspects.hpp
// CHECK-NEXT: info/aspects.def
// CHECK-NEXT: info/aspects_deprecated.def
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: detail/common.hpp
// CHECK-NEXT: vector.hpp
// CHECK-NEXT: detail/named_swizzles_mixin.hpp
// CHECK-NEXT: detail/vector_arith.hpp
// CHECK-NEXT: detail/common.hpp
// CHECK-NEXT: detail/fwd/accessor.hpp
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: detail/memcpy.hpp
// CHECK-NEXT: detail/builtins/common_functions.inc
// CHECK-NEXT: detail/builtins/helper_macros.hpp
// CHECK-NEXT: detail/builtins/geometric_functions.inc
Expand Down
3 changes: 1 addition & 2 deletions sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -97,11 +97,10 @@
// CHECK-NEXT: detail/memcpy.hpp
// CHECK-NEXT: builtins.hpp
// CHECK-NEXT: detail/builtins/builtins.hpp
// CHECK-NEXT: detail/vector_convert.hpp
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: vector.hpp
// CHECK-NEXT: detail/named_swizzles_mixin.hpp
// CHECK-NEXT: detail/vector_arith.hpp
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: detail/builtins/common_functions.inc
// CHECK-NEXT: detail/builtins/helper_macros.hpp
// CHECK-NEXT: detail/builtins/geometric_functions.inc
Expand Down
5 changes: 2 additions & 3 deletions sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,12 +91,11 @@
// CHECK-NEXT: __spirv/spirv_ops.hpp
// CHECK-NEXT: builtins.hpp
// CHECK-NEXT: detail/builtins/builtins.hpp
// CHECK-NEXT: detail/vector_convert.hpp
// CHECK-NEXT: detail/memcpy.hpp
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: vector.hpp
// CHECK-NEXT: detail/named_swizzles_mixin.hpp
// CHECK-NEXT: detail/vector_arith.hpp
// CHECK-NEXT: marray.hpp
// CHECK-NEXT: detail/memcpy.hpp
// CHECK-NEXT: detail/builtins/common_functions.inc
// CHECK-NEXT: detail/builtins/helper_macros.hpp
// CHECK-NEXT: detail/builtins/geometric_functions.inc
Expand Down
Loading
Loading