diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp index d69c0b93ce7c1..9e292c8c78083 100644 --- a/sycl/include/sycl/detail/builtins/builtins.hpp +++ b/sycl/include/sycl/detail/builtins/builtins.hpp @@ -63,12 +63,12 @@ #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,61 @@ 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, + is_multi_ptr_v>> + ? 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, + 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) { + return get_raw_pointer(std::forward(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, + 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) { + return detail::builtin_element_ptr(std::forward(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{}); +} + // 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..91b3b429bf129 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,16 @@ 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..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,16 +10,16 @@ // 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/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..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,16 +11,16 @@ // 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/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..8e3ba08280d87 --- /dev/null +++ b/sycl/test/regression/builtins_multi_ptr_include_order.cpp @@ -0,0 +1,70 @@ +// 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) +// clang-format off +#include +#include +// clang-format on + +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