Skip to content

Commit

Permalink
[SYCL] Remove deprecated raw-pointer math functions overloads (#13238)
Browse files Browse the repository at this point in the history
  • Loading branch information
@aelovikov-intel
aelovikov-intel authored Apr 3, 2024
1 parent f2c5eea commit c3dc8e3
Show file tree
Hide file tree
Showing 3 changed files with 34 additions and 47 deletions.
37 changes: 2 additions & 35 deletions sycl/include/sycl/detail/builtins/math_functions.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -284,22 +284,11 @@ using builtin_last_raw_intptr_t =
change_elements_t<std::conditional_t<is_marray_v<T>, int, int32_t>,
simplify_if_swizzle_t<T>> *;
}
#define BUILTIN_LAST_RAW_INTPTR(NUM_ARGS, NAME) \
template <SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TYPENAME_TYPE)> \
__SYCL_DEPRECATED("SYCL builtin functions with raw pointer arguments have " \
"been deprecated. Please use multi_ptr.") \
detail::builtin_enable_math_allow_scalar_t<SYCL_CONCAT(LESS_ONE(NUM_ARGS), \
_TEMPLATE_TYPE)> \
NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), \
detail::builtin_last_raw_intptr_t<T0> p) { \
return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
}

#ifdef __SYCL_DEVICE_ONLY__
#define BUILTIN_LAST_INTPTR(NUM_ARGS, NAME) \
BUILTIN_LAST_PTR(NUM_ARGS, NAME, builtin_enable_last_intptr_scalar_t, \
builtin_enable_last_intptr_non_scalar_t) \
BUILTIN_LAST_RAW_INTPTR(NUM_ARGS, NAME)
builtin_enable_last_intptr_non_scalar_t)
#else
#define LAST_INT_PTR_DECLARE_SCALAR(NUM_ARGS, NAME, TYPE) \
__SYCL_EXPORT TYPE NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TYPE)(TYPE), \
Expand All @@ -321,8 +310,7 @@ using builtin_last_raw_intptr_t =
} \
} /* namespace detail */ \
BUILTIN_LAST_PTR(NUM_ARGS, NAME, builtin_enable_last_intptr_scalar_t, \
builtin_enable_last_intptr_non_scalar_t) \
BUILTIN_LAST_RAW_INTPTR(NUM_ARGS, NAME)
builtin_enable_last_intptr_non_scalar_t)
#endif

BUILTIN_LAST_INTPTR(TWO_ARGS, frexp)
Expand All @@ -344,13 +332,6 @@ template <typename T0, typename T1> auto fract_impl(T0 &x, T1 &y) {
#endif
BUILTIN_LAST_PTR(TWO_ARGS, fract, builtin_enable_ptr_scalar_t,
builtin_enable_ptr_non_scalar_t)
template <typename T0>
__SYCL_DEPRECATED("SYCL builtin functions with raw pointer arguments have been "
"deprecated. Please use multi_ptr.")
detail::builtin_enable_math_allow_scalar_t<T0> fract(
T0 x, detail::simplify_if_swizzle_t<T0> *y) {
return detail::fract_impl(x, y);
}

#ifndef __SYCL_DEVICE_ONLY__
namespace detail {
Expand All @@ -372,13 +353,6 @@ template <typename T0, typename T1> auto modf_impl(T0 &x, T1 &&y) {
#endif
BUILTIN_LAST_PTR(TWO_ARGS, modf, builtin_enable_ptr_scalar_t,
builtin_enable_ptr_non_scalar_t)
template <typename T0>
__SYCL_DEPRECATED("SYCL builtin functions with raw pointer arguments have been "
"deprecated. Please use multi_ptr.")
detail::builtin_enable_math_allow_scalar_t<T0> modf(
T0 x, detail::simplify_if_swizzle_t<T0> *y) {
return detail::modf_impl(x, y);
}

#undef BUILTIN_LAST_PTR

Expand Down Expand Up @@ -457,13 +431,6 @@ template <typename T0, typename T1> auto sincos_impl(T0 &x, T1 &&y) {
} // namespace detail
BUILTIN_LAST_PTR_COMMON(TWO_ARGS, sincos, builtin_enable_ptr_scalar_t,
builtin_enable_ptr_non_scalar_t)
template <typename T0>
__SYCL_DEPRECATED("SYCL builtin functions with raw pointer arguments have been "
"deprecated. Please use multi_ptr.")
detail::builtin_enable_math_allow_scalar_t<T0> sincos(
T0 x, detail::simplify_if_swizzle_t<T0> *y) {
return detail::sincos_impl(x, y);
}

#undef BUILTIN_LAST_PTR_COMMON
#undef LAST_PTR_SCALAR
Expand Down
34 changes: 26 additions & 8 deletions sycl/test-e2e/USM/math.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -28,8 +28,12 @@ int main() {
void *ptr =
s::malloc_shared(100, myQueue.get_device(), myQueue.get_context());
myQueue.submit([&](s::handler &cgh) {
cgh.single_task<class fractF1UF1>(
[=]() { Buf[0] = s::fract(float{1.5f}, &Buf[1]); });
cgh.single_task<class fractF1UF1>([=]() {
Buf[0] = s::fract(
float{1.5f},
s::address_space_cast<s::access::address_space::global_space,
s::access::decorated::yes>(Buf + 1));
});
});
myQueue.wait();
r = Buf[0];
Expand All @@ -47,7 +51,10 @@ int main() {
sizeof(s::float2) * 2, myQueue.get_device(), myQueue.get_context());
myQueue.submit([&](s::handler &cgh) {
cgh.single_task<class fractF2UF2>([=]() {
Buf[0] = s::fract(s::float2{1.5f, 2.5f}, &Buf[1]);
Buf[0] = s::fract(
s::float2{1.5f, 2.5f},
s::address_space_cast<s::access::address_space::global_space,
s::access::decorated::yes>(Buf + 1));
});
});
myQueue.wait();
Expand All @@ -74,8 +81,12 @@ int main() {
sizeof(int) * 2, myQueue.get_device(), myQueue.get_context());
myQueue.submit([&](s::handler &cgh) {
auto AccR = BufR.get_access<s::access::mode::read_write>(cgh);
cgh.single_task<class lgamma_rF1PI1>(
[=]() { AccR[0] = s::lgamma_r(float{10.f}, BufI); });
cgh.single_task<class lgamma_rF1PI1>([=]() {
AccR[0] = s::lgamma_r(
float{10.f},
s::address_space_cast<s::access::address_space::global_space,
s::access::decorated::yes>(BufI));
});
});
myQueue.wait();
i = *BufI;
Expand All @@ -95,8 +106,12 @@ int main() {
sizeof(int) * 2, myQueue.get_device(), myQueue.get_context());
myQueue.submit([&](s::handler &cgh) {
auto AccR = BufR.get_access<s::access::mode::read_write>(cgh);
cgh.single_task<class lgamma_rF1PI1_neg>(
[=]() { AccR[0] = s::lgamma_r(float{-2.4f}, BufI); });
cgh.single_task<class lgamma_rF1PI1_neg>([=]() {
AccR[0] = s::lgamma_r(
float{-2.4f},
s::address_space_cast<s::access::address_space::global_space,
s::access::decorated::yes>(BufI));
});
});
myQueue.wait();
i = *BufI;
Expand All @@ -117,7 +132,10 @@ int main() {
myQueue.submit([&](s::handler &cgh) {
auto AccR = BufR.get_access<s::access::mode::read_write>(cgh);
cgh.single_task<class lgamma_rF2PF2>([=]() {
AccR[0] = s::lgamma_r(s::float2{10.f, -2.4f}, BufI);
AccR[0] = s::lgamma_r(
s::float2{10.f, -2.4f},
s::address_space_cast<s::access::address_space::global_space,
s::access::decorated::yes>(BufI));
});
});
myQueue.wait();
Expand Down
10 changes: 6 additions & 4 deletions sycl/test/regression/half_ptr_builtins.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,9 +12,11 @@

#include <sycl/sycl.hpp>

int main() {
sycl::half x;
sycl::modf(sycl::half{1.0}, &x);
sycl::sincos(sycl::half{1.0}, &x);
SYCL_EXTERNAL auto
foo(sycl::multi_ptr<sycl::half, sycl::access::address_space::global_space,
sycl::access::decorated::yes>
ptr) {
sycl::modf(sycl::half{1.0}, ptr);
sycl::sincos(sycl::half{1.0}, ptr);
return 0;
}

0 comments on commit c3dc8e3

Please sign in to comment.