Skip to content

Commit f39f1de

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL] Deprecate old sycl_ext_oneapi_group_load_store (intel#15405)
Newer version is under `proposed/`, I'm going to move it to `experimental/` in a separate PR for a cleaner `git diff`.
1 parent c77c419 commit f39f1de

16 files changed

+165
-104
lines changed

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_group_load_store.asciidoc renamed to ‎sycl/doc/extensions/deprecated/sycl_ext_oneapi_group_load_store.asciidoc

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -23,12 +23,8 @@ NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are
2323
trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc.
2424
used by permission by Khronos.
2525

26-
NOTE: This extension is experimental: interfaces are subject to change later.
27-
28-
NOTE: This extension documents functionality that predates SYCL 2020's formal
29-
extension mechanism. Work is underway to align the extension with SYCL 2020;
30-
the documentation in its current state does not reflect the intended long-term
31-
direction of the extension.
26+
NOTE: This extension has been replaced with a new version under the same name
27+
that completely changed the interfaces.
3228

3329
== Notice
3430

‎sycl/include/sycl/sub_group.hpp

Lines changed: 26 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -209,6 +209,7 @@ struct sub_group {
209209
#ifdef __SYCL_DEVICE_ONLY__
210210
// Method for decorated pointer
211211
template <typename CVT, typename T = std::remove_cv_t<CVT>>
212+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
212213
std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value, T>
213214
load(CVT *cv_src) const {
214215
T *src = const_cast<T *>(cv_src);
@@ -219,6 +220,7 @@ struct sub_group {
219220

220221
// Method for raw pointer
221222
template <typename CVT, typename T = std::remove_cv_t<CVT>>
223+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
222224
std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value, T>
223225
load(CVT *cv_src) const {
224226
T *src = const_cast<T *>(cv_src);
@@ -240,6 +242,7 @@ struct sub_group {
240242
}
241243
#else //__SYCL_DEVICE_ONLY__
242244
template <typename CVT, typename T = std::remove_cv_t<CVT>>
245+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
243246
T load(CVT *src) const {
244247
(void)src;
245248
throw sycl::exception(make_error_code(errc::feature_not_supported),
@@ -249,6 +252,7 @@ struct sub_group {
249252

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

270274
template <typename CVT, access::address_space Space,
271275
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
276+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
272277
std::enable_if_t<
273278
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value, T>
274279
load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
@@ -286,6 +291,7 @@ struct sub_group {
286291
#if defined(__NVPTX__) || defined(__AMDGCN__)
287292
template <int N, typename CVT, access::address_space Space,
288293
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
294+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
289295
std::enable_if_t<
290296
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
291297
vec<T, N>>
@@ -301,6 +307,7 @@ struct sub_group {
301307
#else // __NVPTX__ || __AMDGCN__
302308
template <int N, typename CVT, access::address_space Space,
303309
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
310+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
304311
std::enable_if_t<
305312
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
306313
N != 1 && N != 3 && N != 16,
@@ -313,6 +320,7 @@ struct sub_group {
313320

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

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

342351
template <int N, typename CVT, access::address_space Space,
343352
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
353+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
344354
std::enable_if_t<
345355
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
346356
N == 1,
@@ -354,6 +364,7 @@ struct sub_group {
354364
#else // __SYCL_DEVICE_ONLY__
355365
template <int N, typename CVT, access::address_space Space,
356366
access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
367+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
357368
std::enable_if_t<
358369
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
359370
vec<T, N>>
@@ -366,6 +377,7 @@ struct sub_group {
366377

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

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

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

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

451468
template <typename T, access::address_space Space,
452469
access::decorated DecorateAddress>
470+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
453471
std::enable_if_t<
454472
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value>
455473
store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) const {
@@ -467,6 +485,7 @@ struct sub_group {
467485
#if defined(__NVPTX__) || defined(__AMDGCN__)
468486
template <int N, typename T, access::address_space Space,
469487
access::decorated DecorateAddress>
488+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
470489
std::enable_if_t<
471490
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
472491
store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
@@ -477,6 +496,7 @@ struct sub_group {
477496
#else // __NVPTX__ || __AMDGCN__
478497
template <int N, typename T, access::address_space Space,
479498
access::decorated DecorateAddress>
499+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
480500
std::enable_if_t<
481501
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
482502
N != 1 && N != 3 && N != 16>
@@ -486,6 +506,7 @@ struct sub_group {
486506

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

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

507529
template <int N, typename T, access::address_space Space,
508530
access::decorated DecorateAddress>
531+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
509532
std::enable_if_t<
510533
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
511534
N == 16>
@@ -519,6 +542,7 @@ struct sub_group {
519542
#else // __SYCL_DEVICE_ONLY__
520543
template <int N, typename T, access::address_space Space,
521544
access::decorated DecorateAddress>
545+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
522546
std::enable_if_t<
523547
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
524548
store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
@@ -531,6 +555,7 @@ struct sub_group {
531555

532556
template <int N, typename T, access::address_space Space,
533557
access::decorated DecorateAddress>
558+
__SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
534559
std::enable_if_t<
535560
sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value>
536561
store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {

‎sycl/test-e2e/InvokeSimd/Feature/popcnt_emu.cpp

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,7 @@
3636

3737
#include <sycl/detail/core.hpp>
3838
#include <sycl/ext/intel/esimd.hpp>
39+
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
3940
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
4041

4142
#include <functional>
@@ -126,9 +127,9 @@ int main(void) {
126127
} else {
127128
res = id % 2;
128129
}
129-
sg.store(out_accessor.get_multi_ptr<access::decorated::yes>() +
130-
offset,
131-
res);
130+
group_store(sg, res,
131+
out_accessor.get_multi_ptr<access::decorated::yes>() +
132+
offset);
132133
});
133134
});
134135
e.wait();

‎sycl/test-e2e/InvokeSimd/Feature/scale.cpp

Lines changed: 9 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,7 @@
2323

2424
#include <sycl/detail/core.hpp>
2525
#include <sycl/ext/intel/esimd.hpp>
26+
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
2627
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
2728

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

104-
T va = sg.load(
105-
acca.template get_multi_ptr<access::decorated::yes>().get() +
106-
offset);
105+
T va;
106+
group_load(sg,
107+
acca.template get_multi_ptr<access::decorated::yes>() +
108+
offset,
109+
va);
107110
T vc = invoke_simd(sg, SIMD_CALLEE_scale<T>, va, uniform{n});
108-
sg.store(
111+
group_store(
112+
sg, vc,
109113
accc.template get_multi_ptr<access::decorated::yes>().get() +
110-
offset,
111-
vc);
114+
offset);
112115
});
113116
});
114117
e.wait();

‎sycl/test-e2e/InvokeSimd/Feature/void_retval.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,7 @@
2222

2323
#include <sycl/detail/core.hpp>
2424
#include <sycl/ext/intel/esimd.hpp>
25+
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
2526
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
2627

2728
#include <functional>
@@ -102,10 +103,9 @@ int main(void) {
102103

103104
unsigned int offset = g.get_group_id() * g.get_local_range() +
104105
sg.get_group_id() * sg.get_max_local_range();
105-
float va = sg.load(
106-
PA.get_multi_ptr<access::decorated::yes>().get() + offset);
107-
float vb = sg.load(
108-
PB.get_multi_ptr<access::decorated::yes>().get() + offset);
106+
float va, vb;
107+
group_load(sg, PA.get_multi_ptr<access::decorated::yes>().get() + offset, va);
108+
group_load(sg, PB.get_multi_ptr<access::decorated::yes>().get() + offset, vb);
109109
// We need to get a pointer to the starting address of where the
110110
// result of the vector addition should be stored in/written back to
111111
// C. Returns the index (ordinal number) of the work-group to which

‎sycl/test-e2e/InvokeSimd/Regression/call_vadd_1d_spill.cpp

Lines changed: 20 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,7 @@
1717

1818
#include <sycl/detail/core.hpp>
1919
#include <sycl/ext/intel/esimd.hpp>
20+
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
2021
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
2122

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

109110
unsigned int offset = g.get_group_id() * g.get_local_range() +
110111
sg.get_group_id() * sg.get_max_local_range();
111-
float va = sg.load(
112-
PA.get_multi_ptr<access::decorated::yes>().get() + offset);
113-
float vb = sg.load(
114-
PB.get_multi_ptr<access::decorated::yes>().get() + offset);
115-
float vp = sg.load(
116-
PP.get_multi_ptr<access::decorated::yes>().get() + offset);
117-
float vq = sg.load(
118-
PQ.get_multi_ptr<access::decorated::yes>().get() + offset);
119-
float vr = sg.load(
120-
PR.get_multi_ptr<access::decorated::yes>().get() + offset);
121-
float vx = sg.load(
122-
PX.get_multi_ptr<access::decorated::yes>().get() + offset);
123-
float vy = sg.load(
124-
PY.get_multi_ptr<access::decorated::yes>().get() + offset);
125-
float vz = sg.load(
126-
PZ.get_multi_ptr<access::decorated::yes>().get() + offset);
112+
auto Load = [&](auto Acc) {
113+
float res;
114+
group_load(sg,
115+
Acc.template get_multi_ptr<access::decorated::yes>() +
116+
offset,
117+
res);
118+
return res;
119+
};
120+
float va = Load(PA);
121+
float vb = Load(PB);
122+
float vp = Load(PP);
123+
float vq = Load(PQ);
124+
float vr = Load(PR);
125+
float vx = Load(PX);
126+
float vy = Load(PY);
127+
float vz = Load(PZ);
127128

128129
float vc;
129130

@@ -134,8 +135,9 @@ bool test(QueueTY q, float *A, float *B, float *C, float *P, float *Q, float *R,
134135
vc = SPMD_CALLEE_doVadd(va, vb, vx, vy, vx, vy, vx, vy, vx, vy,
135136
vp, vq, vr, vz);
136137
}
137-
sg.store(PC.get_multi_ptr<access::decorated::yes>().get() + offset,
138-
vc);
138+
group_store(sg, vc,
139+
PC.get_multi_ptr<access::decorated::yes>().get() +
140+
offset);
139141
});
140142
});
141143
e.wait();

‎sycl/test-e2e/InvokeSimd/Regression/debug_symbols.cpp

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@
1212

1313
#include <sycl/detail/core.hpp>
1414
#include <sycl/ext/intel/esimd.hpp>
15+
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
1516
#include <sycl/ext/oneapi/experimental/invoke_simd.hpp>
1617

1718
#include <functional>
@@ -84,18 +85,16 @@ int main(void) {
8485

8586
unsigned int offset = g.get_group_id() * g.get_local_range() +
8687
sg.get_group_id() * sg.get_max_local_range();
87-
float va = sg.load(
88-
PA.get_multi_ptr<access::decorated::yes>().get() + offset);
89-
float vb = sg.load(
90-
PB.get_multi_ptr<access::decorated::yes>().get() + offset);
88+
float va, vb;
89+
group_load(sg, PA.get_multi_ptr<access::decorated::yes>().get() + offset, va);
90+
group_load(sg, PB.get_multi_ptr<access::decorated::yes>().get() + offset, vb);
9191

9292
// Invoke SIMD function:
9393
// va values from each work-item are combined into a simd<float,
9494
// VL>. vb values from each work-item are combined into a
9595
// simd<float, VL>.
9696
float vc = invoke_simd(sg, SIMD_CALLEE_doVadd, va, vb);
97-
sg.store(PC.get_multi_ptr<access::decorated::yes>().get() + offset,
98-
vc);
97+
group_store(sg, vc, PC.get_multi_ptr<access::decorated::yes>() + offset);
9998
});
10099
});
101100
e.wait();

0 commit comments

Comments
 (0)