Skip to content

[SYCL] Deprecate old sycl_ext_oneapi_group_load_store #15405

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
merged 3 commits into from
Sep 17, 2024
Merged

[SYCL] Deprecate old sycl_ext_oneapi_group_load_store #15405

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
8d82bf7
[SYCL] Deprecate old `sycl_ext_oneapi_group_load_store`
aelovikov-intel Sep 16, 2024
6de98e2
Fix E2E tests
aelovikov-intel Sep 16, 2024
0ddee44
One more leftover fix in e2e
aelovikov-intel Sep 17, 2024
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
8 changes: 2 additions & 6 deletions ...sycl_ext_oneapi_group_load_store.asciidoc → ...sycl_ext_oneapi_group_load_store.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,12 +23,8 @@ NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are
trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc.
used by permission by Khronos.

NOTE: This extension is experimental: interfaces are subject to change later.

NOTE: This extension documents functionality that predates SYCL 2020's formal
extension mechanism. Work is underway to align the extension with SYCL 2020;
the documentation in its current state does not reflect the intended long-term
direction of the extension.
NOTE: This extension has been replaced with a new version under the same name
that completely changed the interfaces.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In the next PR, can you add a link here to the new spec?


== Notice

Expand Down
27 changes: 26 additions & 1 deletion sycl/include/sycl/sub_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,6 +209,7 @@ struct sub_group {
#ifdef __SYCL_DEVICE_ONLY__
// Method for decorated pointer
template <typename CVT, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value, T>
load(CVT *cv_src) const {
T *src = const_cast<T *>(cv_src);
Expand All @@ -219,6 +220,7 @@ struct sub_group {

// Method for raw pointer
template <typename CVT, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value, T>
load(CVT *cv_src) const {
T *src = const_cast<T *>(cv_src);
Expand All @@ -240,6 +242,7 @@ struct sub_group {
}
#else //__SYCL_DEVICE_ONLY__
template <typename CVT, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
T load(CVT *src) const {
(void)src;
throw sycl::exception(make_error_code(errc::feature_not_supported),
Expand All @@ -249,6 +252,7 @@ struct sub_group {

template <typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value, T>
load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
Expand All @@ -269,6 +273,7 @@ struct sub_group {

template <typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value, T>
load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
Expand All @@ -286,6 +291,7 @@ struct sub_group {
#if defined(__NVPTX__) || defined(__AMDGCN__)
template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
vec<T, N>>
Expand All @@ -301,6 +307,7 @@ struct sub_group {
#else // __NVPTX__ || __AMDGCN__
template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N != 1 && N != 3 && N != 16,
Expand All @@ -313,6 +320,7 @@ struct sub_group {

template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 16,
Expand All @@ -327,6 +335,7 @@ struct sub_group {

template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 3,
Expand All @@ -341,6 +350,7 @@ struct sub_group {

template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 1,
Expand All @@ -354,6 +364,7 @@ struct sub_group {
#else // __SYCL_DEVICE_ONLY__
template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
vec<T, N>>
Expand All @@ -366,6 +377,7 @@ struct sub_group {

template <int N, typename CVT, access::address_space Space,
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value,
vec<T, N>>
Expand All @@ -388,6 +400,7 @@ struct sub_group {
#ifdef __SYCL_DEVICE_ONLY__
// Method for decorated pointer
template <typename T>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value>
store(T *dst, const remove_decoration_t<T> &x) const {
store(sycl::multi_ptr<remove_decoration_t<T>,
Expand All @@ -398,6 +411,7 @@ struct sub_group {

// Method for raw pointer
template <typename T>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value>
store(T *dst, const remove_decoration_t<T> &x) const {

Expand All @@ -421,7 +435,9 @@ struct sub_group {
#endif // __NVPTX__ || __AMDGCN__
}
#else //__SYCL_DEVICE_ONLY__
template <typename T> void store(T *dst, const T &x) const {
template <typename T>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
void store(T *dst, const T &x) const {
(void)dst;
(void)x;
throw sycl::exception(make_error_code(errc::feature_not_supported),
Expand All @@ -431,6 +447,7 @@ struct sub_group {

template <typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) const {
Expand All @@ -450,6 +467,7 @@ struct sub_group {

template <typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value>
store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) const {
Expand All @@ -467,6 +485,7 @@ struct sub_group {
#if defined(__NVPTX__) || defined(__AMDGCN__)
template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
Expand All @@ -477,6 +496,7 @@ struct sub_group {
#else // __NVPTX__ || __AMDGCN__
template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N != 1 && N != 3 && N != 16>
Expand All @@ -486,6 +506,7 @@ struct sub_group {

template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 1>
Expand All @@ -495,6 +516,7 @@ struct sub_group {

template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 3>
Expand All @@ -506,6 +528,7 @@ struct sub_group {

template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 16>
Expand All @@ -519,6 +542,7 @@ struct sub_group {
#else // __SYCL_DEVICE_ONLY__
template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
Expand All @@ -531,6 +555,7 @@ struct sub_group {

template <int N, typename T, access::address_space Space,
access::decorated DecorateAddress>
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
std::enable_if_t<
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value>
store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
Expand Down
7 changes: 4 additions & 3 deletions sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@

#include <sycl/detail/core.hpp>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>

#include <functional>
Expand Down Expand Up @@ -126,9 +127,9 @@ int main(void) {
} else {
res = id % 2;
}
sg.store(out_accessor.get_multi_ptr<access::decorated::yes>() +
offset,
res);
group_store(sg, res,
out_accessor.get_multi_ptr<access::decorated::yes>() +
offset);
});
});
e.wait();
Expand Down
15 changes: 9 additions & 6 deletions sycl/test-e2e/InvokeSimd/Feature/scale.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@

#include <sycl/detail/core.hpp>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>

#include <functional>
Expand Down Expand Up @@ -101,14 +102,16 @@ template <class T, class QueueTY> bool test(QueueTY q) {
unsigned int offset = g.get_group_id() * g.get_local_range() +
sg.get_group_id() * sg.get_max_local_range();

T va = sg.load(
acca.template get_multi_ptr<access::decorated::yes>().get() +
offset);
T va;
group_load(sg,
acca.template get_multi_ptr<access::decorated::yes>() +
offset,
va);
T vc = invoke_simd(sg, SIMD_CALLEE_scale<T>, va, uniform{n});
sg.store(
group_store(
sg, vc,
accc.template get_multi_ptr<access::decorated::yes>().get() +
offset,
vc);
offset);
});
});
e.wait();
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@

#include <sycl/detail/core.hpp>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>

#include <functional>
Expand Down Expand Up @@ -102,10 +103,9 @@ int main(void) {

unsigned int offset = g.get_group_id() * g.get_local_range() +
sg.get_group_id() * sg.get_max_local_range();
float va = sg.load(
PA.get_multi_ptr<access::decorated::yes>().get() + offset);
float vb = sg.load(
PB.get_multi_ptr<access::decorated::yes>().get() + offset);
float va, vb;
group_load(sg, PA.get_multi_ptr<access::decorated::yes>().get() + offset, va);
group_load(sg, PB.get_multi_ptr<access::decorated::yes>().get() + offset, vb);
// We need to get a pointer to the starting address of where the
// result of the vector addition should be stored in/written back to
// C. Returns the index (ordinal number) of the work-group to which
Expand Down
38 changes: 20 additions & 18 deletions sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@

#include <sycl/detail/core.hpp>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>

#include <functional>
Expand Down Expand Up @@ -108,22 +109,22 @@ bool test(QueueTY q, float *A, float *B, float *C, float *P, float *Q, float *R,

unsigned int offset = g.get_group_id() * g.get_local_range() +
sg.get_group_id() * sg.get_max_local_range();
float va = sg.load(
PA.get_multi_ptr<access::decorated::yes>().get() + offset);
float vb = sg.load(
PB.get_multi_ptr<access::decorated::yes>().get() + offset);
float vp = sg.load(
PP.get_multi_ptr<access::decorated::yes>().get() + offset);
float vq = sg.load(
PQ.get_multi_ptr<access::decorated::yes>().get() + offset);
float vr = sg.load(
PR.get_multi_ptr<access::decorated::yes>().get() + offset);
float vx = sg.load(
PX.get_multi_ptr<access::decorated::yes>().get() + offset);
float vy = sg.load(
PY.get_multi_ptr<access::decorated::yes>().get() + offset);
float vz = sg.load(
PZ.get_multi_ptr<access::decorated::yes>().get() + offset);
auto Load = [&](auto Acc) {
float res;
group_load(sg,
Acc.template get_multi_ptr<access::decorated::yes>() +
offset,
res);
return res;
};
float va = Load(PA);
float vb = Load(PB);
float vp = Load(PP);
float vq = Load(PQ);
float vr = Load(PR);
float vx = Load(PX);
float vy = Load(PY);
float vz = Load(PZ);

float vc;

Expand All @@ -134,8 +135,9 @@ bool test(QueueTY q, float *A, float *B, float *C, float *P, float *Q, float *R,
vc = SPMD_CALLEE_doVadd(va, vb, vx, vy, vx, vy, vx, vy, vx, vy,
vp, vq, vr, vz);
}
sg.store(PC.get_multi_ptr<access::decorated::yes>().get() + offset,
vc);
group_store(sg, vc,
PC.get_multi_ptr<access::decorated::yes>().get() +
offset);
});
});
e.wait();
Expand Down
11 changes: 5 additions & 6 deletions sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@

#include <sycl/detail/core.hpp>
#include <sycl/ext/intel/esimd.hpp>
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>

#include <functional>
Expand Down Expand Up @@ -84,18 +85,16 @@ int main(void) {

unsigned int offset = g.get_group_id() * g.get_local_range() +
sg.get_group_id() * sg.get_max_local_range();
float va = sg.load(
PA.get_multi_ptr<access::decorated::yes>().get() + offset);
float vb = sg.load(
PB.get_multi_ptr<access::decorated::yes>().get() + offset);
float va, vb;
group_load(sg, PA.get_multi_ptr<access::decorated::yes>().get() + offset, va);
group_load(sg, PB.get_multi_ptr<access::decorated::yes>().get() + offset, vb);

// Invoke SIMD function:
// va values from each work-item are combined into a simd<float,
// VL>. vb values from each work-item are combined into a
// simd<float, VL>.
float vc = invoke_simd(sg, SIMD_CALLEE_doVadd, va, vb);
sg.store(PC.get_multi_ptr<access::decorated::yes>().get() + offset,
vc);
group_store(sg, vc, PC.get_multi_ptr<access::decorated::yes>() + offset);
});
});
e.wait();
Expand Down
Loading
Loading