From 3578e2dd9363b48339df8543e22cebfe5c6c7a82 Mon Sep 17 00:00:00 2001 From: yes Date: Sun, 5 Apr 2026 14:28:27 -0700 Subject: [PATCH 1/4] Replace direct builtins-side multi_ptr usage with forward-declared helper customization points so no longer needs to include the full definition. Preserve multi_ptr builtin behavior once multi_ptr.hpp is included, and update include-deps / regression coverage. On a TU including only we observe 3.5% reduction for host compilation and 21.7% reduction for device compilation. --- .../include/sycl/detail/builtins/builtins.hpp | 64 +++++++++++++---- .../sycl/detail/builtins/math_functions.inc | 15 ++-- sycl/include/sycl/detail/fwd/multi_ptr.hpp | 14 ++++ sycl/include/sycl/multi_ptr.hpp | 15 ++++ .../sycl_khr_includes_math.hpp.cpp | 4 +- .../sycl_khr_includes_usm.hpp.cpp | 6 +- .../builtins_multi_ptr_include_order.cpp | 68 +++++++++++++++++++ 7 files changed, 156 insertions(+), 30 deletions(-) create mode 100644 sycl/test/regression/builtins_multi_ptr_include_order.cpp diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp index d69c0b93ce7c1..4d0c09bedd6f4 100644 --- a/sycl/include/sycl/detail/builtins/builtins.hpp +++ b/sycl/include/sycl/detail/builtins/builtins.hpp @@ -64,11 +64,11 @@ #pragma once #include +#include #include #include #include #include -#include #include namespace sycl { @@ -83,19 +83,6 @@ template struct use_fast_math : std::false_type {}; #endif template constexpr bool use_fast_math_v = use_fast_math::value; -// Utility trait for getting the decoration of a multi_ptr. -template struct get_multi_ptr_decoration; -template -struct get_multi_ptr_decoration< - multi_ptr> { - static constexpr access::decorated value = DecorateAddress; -}; - -template -constexpr access::decorated get_multi_ptr_decoration_v = - get_multi_ptr_decoration::value; - // Utility trait for checking if a multi_ptr has a "writable" address space, // i.e. global, local, private or generic. template struct has_writeable_addr_space : std::false_type {}; @@ -110,6 +97,55 @@ struct has_writeable_addr_space> template constexpr bool has_writeable_addr_space_v = has_writeable_addr_space::value; +enum class builtin_ptr_kind { raw, multi_ptr }; + +template +using builtin_ptr_kind_tag_t = std::integral_constant< + builtin_ptr_kind, + is_multi_ptr_v>> + ? builtin_ptr_kind::multi_ptr + : builtin_ptr_kind::raw>; + +template +decltype(auto) +builtin_raw_ptr(PtrTy &&Ptr, + std::integral_constant) { + return std::forward(Ptr); +} + +template +auto builtin_raw_ptr(PtrTy &&Ptr, + std::integral_constant) { + return get_raw_pointer(std::forward(Ptr)); +} + +template auto builtin_raw_ptr(PtrTy &&Ptr) { + return builtin_raw_ptr(std::forward(Ptr), + builtin_ptr_kind_tag_t{}); +} + +template +decltype(auto) +builtin_element_ptr(PtrTy &&Ptr, + std::integral_constant) { + return &(*std::forward(Ptr))[0]; +} + +template +auto builtin_element_ptr(PtrTy &&Ptr, + std::integral_constant) { + return detail::builtin_element_ptr(std::forward(Ptr)); +} + +template auto builtin_element_ptr(PtrTy &&Ptr) { + return builtin_element_ptr(std::forward(Ptr), + builtin_ptr_kind_tag_t{}); +} + // Utility trait for changing the element type of a type T. If T is a scalar, // the new type replaces T completely. template diff --git a/sycl/include/sycl/detail/builtins/math_functions.inc b/sycl/include/sycl/detail/builtins/math_functions.inc index 8387fe09e6b69..62f5484517ba6 100644 --- a/sycl/include/sycl/detail/builtins/math_functions.inc +++ b/sycl/include/sycl/detail/builtins/math_functions.inc @@ -223,13 +223,7 @@ auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) { // TODO: Optimize for sizes. Make not to violate ANSI-aliasing rules for the // pointer argument. - auto p0 = [&]() { - if constexpr (is_multi_ptr_v) - return address_space_cast>(&(*p)[0]); - else - return &(*p)[0]; - }(); + auto p0 = builtin_element_ptr(p); constexpr auto N = T0::size(); if constexpr (N <= 16) @@ -314,7 +308,8 @@ using builtin_last_raw_intptr_t = PtrTy p) { \ if constexpr (is_multi_ptr_v) { \ /* TODO: Can't really create multi_ptr on host... */ \ - return NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p.get_raw()); \ + return NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), \ + builtin_raw_ptr(p)); \ } else { \ return builtin_delegate_ptr_impl( \ [](auto... xs) { return NAME##_impl(xs...); }, p, \ @@ -355,7 +350,7 @@ template auto modf_impl(T0 &x, T1 &&y) { if constexpr (is_multi_ptr_v>) { // TODO: Spec needs to be clarified, multi_ptr shouldn't be possible on // host. - return modf_impl(x, y.get_raw()); + return modf_impl(x, builtin_raw_ptr(std::forward(y))); } else { return builtin_delegate_ptr_impl( [](auto x, auto y) { return modf_impl(x, y); }, y, @@ -433,7 +428,7 @@ template auto sincos_impl(T0 &x, T1 &&y) { if constexpr (is_multi_ptr_v>) { // TODO: Spec needs to be clarified, multi_ptr shouldn't be possible on // host. - return sincos_impl(x, y.get_raw()); + return sincos_impl(x, builtin_raw_ptr(std::forward(y))); } else { return builtin_delegate_ptr_impl( [](auto... xs) { return sincos_impl(xs...); }, y, diff --git a/sycl/include/sycl/detail/fwd/multi_ptr.hpp b/sycl/include/sycl/detail/fwd/multi_ptr.hpp index c3718463dc070..202116d8c7999 100644 --- a/sycl/include/sycl/detail/fwd/multi_ptr.hpp +++ b/sycl/include/sycl/detail/fwd/multi_ptr.hpp @@ -10,6 +10,8 @@ #include +#include + namespace sycl { inline namespace _V1 { // Forward declaration @@ -20,5 +22,17 @@ template multi_ptr address_space_cast(ElementType *pointer); + +namespace detail { +template +std::add_pointer_t +get_raw_pointer(multi_ptr Ptr); + +template +auto builtin_element_ptr( + multi_ptr Ptr); +} // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp index 3106bd3cc1a6c..ade21a5a6cd9b 100644 --- a/sycl/include/sycl/multi_ptr.hpp +++ b/sycl/include/sycl/multi_ptr.hpp @@ -1377,6 +1377,21 @@ address_space_cast(ElementType *pointer) { pointer)); } +namespace detail { +template +std::add_pointer_t +get_raw_pointer(multi_ptr Ptr) { + return Ptr.get_raw(); +} + +template +auto builtin_element_ptr(multi_ptr Ptr) { + return address_space_cast(&(*Ptr)[0]); +} +} // namespace detail + template < typename ElementType, access::address_space Space, access::decorated DecorateAddress = access::decorated::legacy, diff --git a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp index 6b0b4af4e44f1..7ff05121b1979 100644 --- a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp @@ -17,9 +17,9 @@ // CHECK-NEXT: detail/export.hpp // CHECK-NEXT: memory_enums.hpp // CHECK-NEXT: __spirv/spirv_vars.hpp +// CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: detail/type_traits.hpp // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp -// CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: detail/vector_convert.hpp // CHECK-NEXT: detail/generic_type_traits.hpp // CHECK-NEXT: aliases.hpp @@ -37,8 +37,6 @@ // CHECK-NEXT: detail/common.hpp // CHECK-NEXT: detail/fwd/accessor.hpp // CHECK-NEXT: marray.hpp -// CHECK-NEXT: multi_ptr.hpp -// CHECK-NEXT: detail/address_space_cast.hpp // CHECK-NEXT: detail/builtins/common_functions.inc // CHECK-NEXT: detail/builtins/helper_macros.hpp // CHECK-NEXT: detail/builtins/geometric_functions.inc diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp index e26c64af65dfb..c9dce9a912558 100644 --- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp @@ -18,9 +18,9 @@ // CHECK-NEXT: detail/export.hpp // CHECK-NEXT: memory_enums.hpp // CHECK-NEXT: __spirv/spirv_vars.hpp +// CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: detail/type_traits.hpp // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp -// CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: detail/vector_convert.hpp // CHECK-NEXT: detail/generic_type_traits.hpp // CHECK-NEXT: aliases.hpp @@ -38,8 +38,6 @@ // CHECK-NEXT: detail/common.hpp // CHECK-NEXT: detail/fwd/accessor.hpp // CHECK-NEXT: marray.hpp -// CHECK-NEXT: multi_ptr.hpp -// CHECK-NEXT: detail/address_space_cast.hpp // CHECK-NEXT: detail/builtins/common_functions.inc // CHECK-NEXT: detail/builtins/helper_macros.hpp // CHECK-NEXT: detail/builtins/geometric_functions.inc @@ -102,6 +100,8 @@ // CHECK-NEXT: ext/oneapi/accessor_property_list.hpp // CHECK-NEXT: detail/accessor_iterator.hpp // CHECK-NEXT: detail/handler_proxy.hpp +// CHECK-NEXT: multi_ptr.hpp +// CHECK-NEXT: detail/address_space_cast.hpp // CHECK-NEXT: pointers.hpp // CHECK-NEXT: properties/accessor_properties.hpp // CHECK-NEXT: properties/runtime_accessor_properties.def diff --git a/sycl/test/regression/builtins_multi_ptr_include_order.cpp b/sycl/test/regression/builtins_multi_ptr_include_order.cpp new file mode 100644 index 0000000000000..b8f2033958e76 --- /dev/null +++ b/sycl/test/regression/builtins_multi_ptr_include_order.cpp @@ -0,0 +1,68 @@ +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations %s -DTEST_BUILTINS_ONLY +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations %s -DTEST_BUILTINS_FIRST +// RUN: %clangxx -fsycl -fsyntax-only -Wno-deprecated-declarations %s -DTEST_MULTI_PTR_FIRST + +// Regression coverage for builtins/multi_ptr decoupling. +// We want to preserve these behaviors: +// 1. compiles without including . +// 2. Including builtins before multi_ptr still allows later multi_ptr +// instantiation for scalar pointer builtins. +// 3. Including builtins before multi_ptr still allows later multi_ptr +// instantiation for vector pointer builtins. +// 4. Including multi_ptr before builtins also works for those builtin calls. + +#if defined(TEST_BUILTINS_ONLY) +#include + +int main() { + auto Value = sycl::fmin(1.0f, 2.0f); + (void)Value; + return 0; +} + +#elif defined(TEST_BUILTINS_FIRST) +#include +#include + +SYCL_EXTERNAL void testScalar( + sycl::multi_ptr + Ptr) { + (void)sycl::modf(1.0f, Ptr); + (void)sycl::sincos(1.0f, Ptr); +} + +SYCL_EXTERNAL void testVector( + sycl::multi_ptr, + sycl::access::address_space::global_space, + sycl::access::decorated::no> + Ptr) { + sycl::vec Value{1.0f, 2.0f}; + (void)sycl::fract(Value, Ptr); +} + +int main() { return 0; } + +#elif defined(TEST_MULTI_PTR_FIRST) +#include +#include + +SYCL_EXTERNAL void testScalar( + sycl::multi_ptr + Ptr) { + (void)sycl::modf(1.0f, Ptr); + (void)sycl::sincos(1.0f, Ptr); +} + +SYCL_EXTERNAL void testVector( + sycl::multi_ptr, + sycl::access::address_space::global_space, + sycl::access::decorated::no> + Ptr) { + sycl::vec Value{1.0f, 2.0f}; + (void)sycl::fract(Value, Ptr); +} + +int main() { return 0; } +#endif \ No newline at end of file From 131bd24d161ce5115ae729179e6fb99366359071 Mon Sep 17 00:00:00 2001 From: yes Date: Mon, 6 Apr 2026 09:58:18 -0700 Subject: [PATCH 2/4] Add comments --- sycl/include/sycl/detail/builtins/builtins.hpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp index 4d0c09bedd6f4..3908f397b1396 100644 --- a/sycl/include/sycl/detail/builtins/builtins.hpp +++ b/sycl/include/sycl/detail/builtins/builtins.hpp @@ -97,8 +97,10 @@ struct has_writeable_addr_space> template constexpr bool has_writeable_addr_space_v = has_writeable_addr_space::value; +// Classification of pointer-like types used by builtin pointer helpers. enum class builtin_ptr_kind { raw, multi_ptr }; +// Maps a pointer-like type to the corresponding builtin_ptr_kind tag. template using builtin_ptr_kind_tag_t = std::integral_constant< builtin_ptr_kind, @@ -106,6 +108,7 @@ using builtin_ptr_kind_tag_t = std::integral_constant< ? builtin_ptr_kind::multi_ptr : builtin_ptr_kind::raw>; +// Returns Ptr unchanged for raw pointer-like types. template decltype(auto) builtin_raw_ptr(PtrTy &&Ptr, @@ -114,6 +117,7 @@ builtin_raw_ptr(PtrTy &&Ptr, return std::forward(Ptr); } +// Extracts the underlying raw pointer from a multi_ptr. template auto builtin_raw_ptr(PtrTy &&Ptr, std::integral_constant(Ptr)); } +// Returns a raw pointer representation for raw pointers and multi_ptrs. template auto builtin_raw_ptr(PtrTy &&Ptr) { return builtin_raw_ptr(std::forward(Ptr), builtin_ptr_kind_tag_t{}); } +// Returns a pointer to the first element for raw pointer-like types. template decltype(auto) builtin_element_ptr(PtrTy &&Ptr, @@ -134,6 +140,7 @@ builtin_element_ptr(PtrTy &&Ptr, return &(*std::forward(Ptr))[0]; } +// Returns a pointer to the first element while preserving multi_ptr semantics. template auto builtin_element_ptr(PtrTy &&Ptr, std::integral_constant(Ptr)); } +// Returns an element pointer for raw pointers and multi_ptrs. template auto builtin_element_ptr(PtrTy &&Ptr) { return builtin_element_ptr(std::forward(Ptr), builtin_ptr_kind_tag_t{}); From 7e9c782781c9248d4996bcf5295993332df76820 Mon Sep 17 00:00:00 2001 From: yes Date: Mon, 6 Apr 2026 12:27:14 -0700 Subject: [PATCH 3/4] Fix clang-format --- .../include/sycl/detail/builtins/builtins.hpp | 28 +++++++------- sycl/include/sycl/detail/fwd/multi_ptr.hpp | 3 +- .../builtins_multi_ptr_include_order.cpp | 38 ++++++++++--------- 3 files changed, 34 insertions(+), 35 deletions(-) diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp index 3908f397b1396..9e292c8c78083 100644 --- a/sycl/include/sycl/detail/builtins/builtins.hpp +++ b/sycl/include/sycl/detail/builtins/builtins.hpp @@ -63,8 +63,8 @@ #pragma once -#include #include +#include #include #include #include @@ -110,18 +110,17 @@ using builtin_ptr_kind_tag_t = std::integral_constant< // Returns Ptr unchanged for raw pointer-like types. template -decltype(auto) -builtin_raw_ptr(PtrTy &&Ptr, - std::integral_constant) { +decltype(auto) builtin_raw_ptr( + PtrTy &&Ptr, + std::integral_constant) { return std::forward(Ptr); } // Extracts the underlying raw pointer from a multi_ptr. template -auto builtin_raw_ptr(PtrTy &&Ptr, - std::integral_constant) { +auto builtin_raw_ptr( + PtrTy &&Ptr, + std::integral_constant) { return get_raw_pointer(std::forward(Ptr)); } @@ -133,18 +132,17 @@ template auto builtin_raw_ptr(PtrTy &&Ptr) { // Returns a pointer to the first element for raw pointer-like types. template -decltype(auto) -builtin_element_ptr(PtrTy &&Ptr, - std::integral_constant) { +decltype(auto) builtin_element_ptr( + PtrTy &&Ptr, + std::integral_constant) { return &(*std::forward(Ptr))[0]; } // Returns a pointer to the first element while preserving multi_ptr semantics. template -auto builtin_element_ptr(PtrTy &&Ptr, - std::integral_constant) { +auto builtin_element_ptr( + PtrTy &&Ptr, + std::integral_constant) { return detail::builtin_element_ptr(std::forward(Ptr)); } diff --git a/sycl/include/sycl/detail/fwd/multi_ptr.hpp b/sycl/include/sycl/detail/fwd/multi_ptr.hpp index 202116d8c7999..91b3b429bf129 100644 --- a/sycl/include/sycl/detail/fwd/multi_ptr.hpp +++ b/sycl/include/sycl/detail/fwd/multi_ptr.hpp @@ -31,8 +31,7 @@ get_raw_pointer(multi_ptr Ptr); template -auto builtin_element_ptr( - multi_ptr Ptr); +auto builtin_element_ptr(multi_ptr Ptr); } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/test/regression/builtins_multi_ptr_include_order.cpp b/sycl/test/regression/builtins_multi_ptr_include_order.cpp index b8f2033958e76..8e3ba08280d87 100644 --- a/sycl/test/regression/builtins_multi_ptr_include_order.cpp +++ b/sycl/test/regression/builtins_multi_ptr_include_order.cpp @@ -24,19 +24,19 @@ int main() { #include #include -SYCL_EXTERNAL void testScalar( - sycl::multi_ptr - Ptr) { +SYCL_EXTERNAL void +testScalar(sycl::multi_ptr + Ptr) { (void)sycl::modf(1.0f, Ptr); (void)sycl::sincos(1.0f, Ptr); } -SYCL_EXTERNAL void testVector( - sycl::multi_ptr, - sycl::access::address_space::global_space, - sycl::access::decorated::no> - Ptr) { +SYCL_EXTERNAL void +testVector(sycl::multi_ptr, + sycl::access::address_space::global_space, + sycl::access::decorated::no> + Ptr) { sycl::vec Value{1.0f, 2.0f}; (void)sycl::fract(Value, Ptr); } @@ -44,22 +44,24 @@ SYCL_EXTERNAL void testVector( int main() { return 0; } #elif defined(TEST_MULTI_PTR_FIRST) +// clang-format off #include #include +// clang-format on -SYCL_EXTERNAL void testScalar( - sycl::multi_ptr - Ptr) { +SYCL_EXTERNAL void +testScalar(sycl::multi_ptr + Ptr) { (void)sycl::modf(1.0f, Ptr); (void)sycl::sincos(1.0f, Ptr); } -SYCL_EXTERNAL void testVector( - sycl::multi_ptr, - sycl::access::address_space::global_space, - sycl::access::decorated::no> - Ptr) { +SYCL_EXTERNAL void +testVector(sycl::multi_ptr, + sycl::access::address_space::global_space, + sycl::access::decorated::no> + Ptr) { sycl::vec Value{1.0f, 2.0f}; (void)sycl::fract(Value, Ptr); } From 2c5fd432362df2f0b13035a6bec54c3ff8f675a0 Mon Sep 17 00:00:00 2001 From: y Date: Mon, 6 Apr 2026 15:09:43 -0700 Subject: [PATCH 4/4] Fix ordering --- sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp | 4 ++-- sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp index 7ff05121b1979..bda8ef3f0adb5 100644 --- a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp @@ -10,14 +10,14 @@ // CHECK-NEXT: feature_test.hpp // CHECK-NEXT: builtins.hpp // CHECK-NEXT: detail/builtins/builtins.hpp +// CHECK-NEXT: detail/fwd/multi_ptr.hpp +// CHECK-NEXT: access/access.hpp // CHECK-NEXT: detail/helpers.hpp // CHECK-NEXT: __spirv/spirv_types.hpp // CHECK-NEXT: detail/defines.hpp -// CHECK-NEXT: access/access.hpp // CHECK-NEXT: detail/export.hpp // CHECK-NEXT: memory_enums.hpp // CHECK-NEXT: __spirv/spirv_vars.hpp -// CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: detail/type_traits.hpp // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp // CHECK-NEXT: detail/vector_convert.hpp diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp index c9dce9a912558..79f0cfcea9a3a 100644 --- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp @@ -11,14 +11,14 @@ // CHECK-NEXT: usm.hpp // CHECK-NEXT: builtins.hpp // CHECK-NEXT: detail/builtins/builtins.hpp +// CHECK-NEXT: detail/fwd/multi_ptr.hpp +// CHECK-NEXT: access/access.hpp // CHECK-NEXT: detail/helpers.hpp // CHECK-NEXT: __spirv/spirv_types.hpp // CHECK-NEXT: detail/defines.hpp -// CHECK-NEXT: access/access.hpp // CHECK-NEXT: detail/export.hpp // CHECK-NEXT: memory_enums.hpp // CHECK-NEXT: __spirv/spirv_vars.hpp -// CHECK-NEXT: detail/fwd/multi_ptr.hpp // CHECK-NEXT: detail/type_traits.hpp // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp // CHECK-NEXT: detail/vector_convert.hpp