Skip to content

Commit 52edb37

Browse files
mmoadelimmoadeli
authored
Implement multi_ptr default to be legacy to avoid code break with SYCL 1.2.1 (#10174)
* This PR covers updates in [Default multi_ptr to legacy to avoid code break with SYCL 1.2.1](KhronosGroup/SYCL-Docs#432) * Declare access::decorated::legacy as deprecated. * Revert `get_pointer` interface for device, host and local_accessor back to 1.2.1 version. * `enable_if` `async_work_group_copy` in `group.hpp` and `nd_item.hpp` to ensure Dst and Src(const / non-const) have the same type. * Add multi_ptr::get_raw and multi_ptr::get_decorated member functions to the access::decorated::legacy specialization of multi_ptr.
1 parent ad88048 commit 52edb37

File tree

15 files changed

+104
-41
lines changed

15 files changed

+104
-41
lines changed

‎sycl/include/sycl/access/access.hpp

Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -56,7 +56,12 @@ enum class address_space : int {
5656
generic_space = 6, // TODO generic_space address space is not supported yet
5757
};
5858

59-
enum class decorated : int { no = 0, yes = 1, legacy = 2 };
59+
enum class decorated : int {
60+
no = 0,
61+
yes = 1,
62+
legacy __SYCL2020_DEPRECATED("sycl::access::decorated::legacy "
63+
"is deprecated since SYCL 2020") = 2
64+
};
6065
} // namespace access
6166

6267
using access::target;

‎sycl/include/sycl/accessor.hpp

Lines changed: 16 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2292,12 +2292,22 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
22922292
template <access::target AccessTarget_ = AccessTarget,
22932293
typename = std::enable_if_t<
22942294
(AccessTarget_ == access::target::host_buffer) ||
2295-
(AccessTarget_ == access::target::host_task) ||
2296-
(AccessTarget_ == access::target::device)>>
2295+
(AccessTarget_ == access::target::host_task)>>
22972296
std::add_pointer_t<value_type> get_pointer() const noexcept {
22982297
return getPointerAdjusted();
22992298
}
23002299

2300+
template <
2301+
access::target AccessTarget_ = AccessTarget,
2302+
typename = std::enable_if_t<(AccessTarget_ == access::target::device)>>
2303+
__SYCL2020_DEPRECATED(
2304+
"accessor::get_pointer() is deprecated, please use get_multi_ptr()")
2305+
global_ptr<DataT> get_pointer() const noexcept {
2306+
return global_ptr<DataT>(
2307+
const_cast<typename detail::DecoratedType<DataT, AS>::type *>(
2308+
getPointerAdjusted()));
2309+
}
2310+
23012311
template <access::target AccessTarget_ = AccessTarget,
23022312
typename = std::enable_if_t<AccessTarget_ ==
23032313
access::target::constant_buffer>>
@@ -3064,8 +3074,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
30643074
return const_reverse_iterator(begin());
30653075
}
30663076

3067-
std::add_pointer_t<value_type> get_pointer() const noexcept {
3068-
return std::add_pointer_t<value_type>(local_acc::getQualifiedPtr());
3077+
__SYCL2020_DEPRECATED(
3078+
"local_accessor::get_pointer() is deprecated, please use get_multi_ptr()")
3079+
local_ptr<DataT> get_pointer() const noexcept {
3080+
return local_ptr<DataT>(local_acc::getQualifiedPtr());
30693081
}
30703082

30713083
template <access::decorated IsDecorated>

‎sycl/include/sycl/ext/intel/esimd/detail/sycl_util.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -82,7 +82,7 @@ using EnableIfAccessor = std::enable_if_t<
8282
template <typename T, int Dimensions>
8383
__ESIMD_API uint32_t localAccessorToOffset(local_accessor<T, Dimensions> acc) {
8484
return static_cast<uint32_t>(
85-
reinterpret_cast<std::uintptr_t>(acc.get_pointer()));
85+
reinterpret_cast<std::uintptr_t>(acc.get_pointer().get()));
8686
}
8787

8888
} // namespace ext::intel::esimd::detail

‎sycl/include/sycl/ext/intel/esimd/detail/util.hpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -189,7 +189,8 @@ auto accessorToPointer(AccessorTy Acc, OffsetTy Offset = 0) {
189189
using QualTPtrType =
190190
std::conditional_t<std::is_const_v<typename AccessorTy::value_type>,
191191
const T *, T *>;
192-
auto BytePtr = reinterpret_cast<QualCharPtrType>(Acc.get_pointer()) + Offset;
192+
auto BytePtr =
193+
reinterpret_cast<QualCharPtrType>(Acc.get_pointer().get()) + Offset;
193194
return reinterpret_cast<QualTPtrType>(BytePtr);
194195
}
195196
#endif // __ESIMD_FORCE_STATELESS_MEM

‎sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -729,7 +729,7 @@ lsc_gather(AccessorTy acc,
729729
__ESIMD_NS::simd_mask<N> pred = 1) {
730730
#ifdef __ESIMD_FORCE_STATELESS_MEM
731731
return lsc_gather<T, NElts, DS, L1H, L3H>(
732-
reinterpret_cast<T *>(acc.get_pointer()), offsets, pred);
732+
reinterpret_cast<T *>(acc.get_pointer().get()), offsets, pred);
733733
#else
734734
detail::check_lsc_vector_size<NElts>();
735735
detail::check_lsc_data_size<T, DS>();
@@ -803,7 +803,8 @@ lsc_gather(AccessorTy acc,
803803
__ESIMD_NS::simd<T, N * NElts> old_values) {
804804
#ifdef __ESIMD_FORCE_STATELESS_MEM
805805
return lsc_gather<T, NElts, DS, L1H, L3H>(
806-
reinterpret_cast<T *>(acc.get_pointer()), offsets, pred, old_values);
806+
reinterpret_cast<T *>(acc.get_pointer().get()), offsets, pred,
807+
old_values);
807808

808809
#else
809810
detail::check_lsc_vector_size<NElts>();

‎sycl/include/sycl/group.hpp

Lines changed: 10 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -526,9 +526,11 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
526526
/// Permitted types for DestDataT are all scalar and vector types. SrcDataT
527527
/// must be either the same as DestDataT or const DestDataT.
528528
template <typename DestDataT, typename SrcDataT>
529-
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
530-
decorated_global_ptr<SrcDataT> src,
531-
size_t numElements) const {
529+
typename std::enable_if_t<
530+
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
531+
async_work_group_copy(decorated_local_ptr<DestDataT> dest,
532+
decorated_global_ptr<SrcDataT> src,
533+
size_t numElements) const {
532534
return async_work_group_copy(dest, src, numElements, 1);
533535
}
534536

@@ -539,9 +541,11 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
539541
/// Permitted types for DestDataT are all scalar and vector types. SrcDataT
540542
/// must be either the same as DestDataT or const DestDataT.
541543
template <typename DestDataT, typename SrcDataT>
542-
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
543-
decorated_local_ptr<SrcDataT> src,
544-
size_t numElements) const {
544+
typename std::enable_if_t<
545+
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
546+
async_work_group_copy(decorated_global_ptr<DestDataT> dest,
547+
decorated_local_ptr<SrcDataT> src,
548+
size_t numElements) const {
545549
return async_work_group_copy(dest, src, numElements, 1);
546550
}
547551

‎sycl/include/sycl/multi_ptr.hpp

Lines changed: 17 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -662,6 +662,7 @@ class multi_ptr<void, Space, DecorateAddress> {
662662
template <typename ElementType, access::address_space Space>
663663
class multi_ptr<ElementType, Space, access::decorated::legacy> {
664664
public:
665+
using value_type = ElementType;
665666
using element_type =
666667
std::conditional_t<std::is_same_v<ElementType, half>,
667668
sycl::detail::half_impl::BIsRepresentationT,
@@ -773,9 +774,8 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
773774
Space == access::address_space::ext_intel_global_device_space)>>
774775
multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
775776
isPlaceholder, PropertyListT>
776-
Accessor) {
777-
m_Pointer = detail::cast_AS<pointer_t>(Accessor.get_pointer());
778-
}
777+
Accessor)
778+
: multi_ptr(detail::cast_AS<pointer_t>(Accessor.get_pointer().get())) {}
779779

780780
// Only if Space == local_space || generic_space
781781
template <
@@ -891,6 +891,10 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
891891

892892
// Returns the underlying OpenCL C pointer
893893
pointer_t get() const { return m_Pointer; }
894+
pointer_t get_decorated() const { return m_Pointer; }
895+
std::add_pointer_t<element_type> get_raw() const {
896+
return reinterpret_cast<std::add_pointer_t<element_type>>(get());
897+
}
894898

895899
// Implicit conversion to the underlying pointer type
896900
operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); }
@@ -1003,6 +1007,7 @@ class multi_ptr<ElementType, Space, access::decorated::legacy> {
10031007
template <access::address_space Space>
10041008
class multi_ptr<void, Space, access::decorated::legacy> {
10051009
public:
1010+
using value_type = void;
10061011
using element_type = void;
10071012
using difference_type = std::ptrdiff_t;
10081013

@@ -1114,6 +1119,10 @@ class multi_ptr<void, Space, access::decorated::legacy> {
11141119
using ReturnPtr = detail::const_if_const_AS<Space, void> *;
11151120
// Returns the underlying OpenCL C pointer
11161121
pointer_t get() const { return m_Pointer; }
1122+
pointer_t get_decorated() const { return m_Pointer; }
1123+
std::add_pointer_t<element_type> get_raw() const {
1124+
return reinterpret_cast<std::add_pointer_t<element_type>>(get());
1125+
}
11171126

11181127
// Implicit conversion to the underlying pointer type
11191128
operator ReturnPtr() const { return reinterpret_cast<ReturnPtr>(m_Pointer); };
@@ -1144,6 +1153,7 @@ class multi_ptr<void, Space, access::decorated::legacy> {
11441153
template <access::address_space Space>
11451154
class multi_ptr<const void, Space, access::decorated::legacy> {
11461155
public:
1156+
using value_type = const void;
11471157
using element_type = const void;
11481158
using difference_type = std::ptrdiff_t;
11491159

@@ -1256,6 +1266,10 @@ class multi_ptr<const void, Space, access::decorated::legacy> {
12561266

12571267
// Returns the underlying OpenCL C pointer
12581268
pointer_t get() const { return m_Pointer; }
1269+
pointer_t get_decorated() const { return m_Pointer; }
1270+
std::add_pointer_t<element_type> get_raw() const {
1271+
return reinterpret_cast<std::add_pointer_t<element_type>>(get());
1272+
}
12591273

12601274
// Implicit conversion to the underlying pointer type
12611275
operator const void *() const {

‎sycl/include/sycl/nd_item.hpp

Lines changed: 20 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -170,33 +170,39 @@ template <int Dimensions = 1> class nd_item {
170170
}
171171

172172
template <typename DestDataT, typename SrcDataT>
173-
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
174-
decorated_global_ptr<SrcDataT> src,
175-
size_t numElements) const {
173+
typename std::enable_if_t<
174+
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
175+
async_work_group_copy(decorated_local_ptr<DestDataT> dest,
176+
decorated_global_ptr<SrcDataT> src,
177+
size_t numElements) const {
176178
return Group.async_work_group_copy(dest, src, numElements);
177179
}
178180

179181
template <typename DestDataT, typename SrcDataT>
180-
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
181-
decorated_local_ptr<SrcDataT> src,
182-
size_t numElements) const {
182+
typename std::enable_if_t<
183+
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
184+
async_work_group_copy(decorated_global_ptr<DestDataT> dest,
185+
decorated_local_ptr<SrcDataT> src,
186+
size_t numElements) const {
183187
return Group.async_work_group_copy(dest, src, numElements);
184188
}
185189

186190
template <typename DestDataT, typename SrcDataT>
187-
device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
188-
decorated_global_ptr<SrcDataT> src,
189-
size_t numElements,
190-
size_t srcStride) const {
191+
typename std::enable_if_t<
192+
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
193+
async_work_group_copy(decorated_local_ptr<DestDataT> dest,
194+
decorated_global_ptr<SrcDataT> src, size_t numElements,
195+
size_t srcStride) const {
191196

192197
return Group.async_work_group_copy(dest, src, numElements, srcStride);
193198
}
194199

195200
template <typename DestDataT, typename SrcDataT>
196-
device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
197-
decorated_local_ptr<SrcDataT> src,
198-
size_t numElements,
199-
size_t destStride) const {
201+
typename std::enable_if_t<
202+
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
203+
async_work_group_copy(decorated_global_ptr<DestDataT> dest,
204+
decorated_local_ptr<SrcDataT> src, size_t numElements,
205+
size_t destStride) const {
200206
return Group.async_work_group_copy(dest, src, numElements, destStride);
201207
}
202208

‎sycl/test-e2e/ESIMD/accessor_local.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -53,7 +53,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
5353
uint32_t GID = Item.get_global_id(0);
5454
uint32_t LID = Item.get_local_id(0);
5555
uint32_t LocalAccOffset = static_cast<uint32_t>(
56-
reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer()));
56+
reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer().get()));
5757
if constexpr (TestSubscript) {
5858
for (int I = 0; I < VL; I++)
5959
LocalAcc[LID * VL + I] = GID * 100 + I;

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -45,7 +45,7 @@ ESIMD_INLINE void slm_load_store_test(
4545

4646
uint32_t LocalAccOffset =
4747
static_cast<uint32_t>(
48-
reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer())) +
48+
reinterpret_cast<std::uintptr_t>(LocalAcc.get_pointer().get())) +
4949
LAByteOffset;
5050
auto Local1 = esimd::slm_block_load<dtype, VL>(LocalAccOffset);
5151
auto Local2 = esimd::slm_block_load<dtype, VL>(LocalAccOffset +

‎sycl/test/basic_tests/accessor/accessor_get_pointer.cpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -27,8 +27,6 @@ void test_get_multi_ptr(handler &cgh, buffer<int, size> &buffer) {
2727
auto device_acc_ptr = device_acc.get_pointer();
2828
static_assert(std::is_same_v<decltype(acc_ptr), std::add_pointer_t<int>>);
2929
static_assert(std::is_same_v<decltype(target_local_ptr), local_ptr<int>>);
30-
static_assert(
31-
std::is_same_v<decltype(local_pointer), std::add_pointer_t<int>>);
32-
static_assert(
33-
std::is_same_v<decltype(device_acc_ptr), std::add_pointer_t<int>>);
30+
static_assert(std::is_same_v<decltype(local_pointer), local_ptr<int>>);
31+
static_assert(std::is_same_v<decltype(device_acc_ptr), global_ptr<int>>);
3432
}

‎sycl/test/esimd/esimd_verify.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
using namespace sycl;
1010
using namespace sycl::ext::intel::esimd;
1111

12-
// CHECK-NEGATIVE-DAG: error: function 'int* sycl::_V1::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
12+
// CHECK-NEGATIVE-DAG: error: function 'sycl::_V1::multi_ptr<{{.+}}> sycl::_V1::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
1313
// CHECK-NEGATIVE-DAG: error: function '{{.+}} sycl::_V1::accessor<{{.+}}>::operator[]<{{.+}}>({{.+}}) const' is not supported in ESIMD context
1414
// CHECK-NEGATIVE-DAG: error: function '{{.+}}combine(int const&)' is not supported in ESIMD context
1515

‎sycl/test/extensions/usm_pointers_aliases.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify %s -o %t.out
1+
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify %s -o %t.out -Wno-deprecated-declarations
22
// expected-no-diagnostics
33

44
#include <sycl/sycl.hpp>

‎sycl/test/multi_ptr/ctad.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -49,8 +49,11 @@ int main() {
4949
sycl::access::decorated::no>;
5050
using constMPtr = sycl::multi_ptr<int, address_space::constant_space,
5151
sycl::access::decorated::legacy>;
52+
using constDefaultMPtr = sycl::multi_ptr<int, address_space::constant_space>;
5253
using localMPtr = sycl::multi_ptr<int, address_space::local_space,
5354
sycl::access::decorated::no>;
55+
using legacyMPtr = sycl::multi_ptr<int, address_space::global_space,
56+
sycl::access::decorated::legacy>;
5457
static_assert(std::is_same<rwDeviceCTAD, deviceMPtr>::value);
5558
static_assert(std::is_same<rwDeviceCTAD, globlMPtr>::value);
5659
static_assert(std::is_same<rwGloblCTAD, globlMPtr>::value);
@@ -63,6 +66,15 @@ int main() {
6366
static_assert(std::is_same<constCTAD, constMPtr>::value);
6467
static_assert(std::is_same<localCTAD, localMPtr>::value);
6568
static_assert(std::is_same<localCTADDep, localMPtr>::value);
69+
static_assert(std::is_same<constMPtr, constDefaultMPtr>::value);
70+
71+
legacyMPtr LegacytMultiPtr;
72+
static_assert(
73+
std::is_same_v<
74+
decltype(LegacytMultiPtr.get_decorated()),
75+
typename sycl::multi_ptr<int, address_space::global_space,
76+
sycl::access::decorated::yes>::pointer>);
77+
static_assert(std::is_same_v<decltype(LegacytMultiPtr.get_raw()), int *>);
6678

6779
globlMPtr non_const_multi_ptr;
6880
auto constTypeMultiPtr = constTypeMPtr(non_const_multi_ptr);

‎sycl/test/warnings/sycl_2020_deprecations.cpp

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -303,21 +303,31 @@ int main() {
303303
[=](sycl::nd_item<1> Idx) {
304304
int PrivateVal = 0;
305305

306+
// expected-warning@+6{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
307+
// expected-warning@+8{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
308+
// expected-warning@+8{{'get_pointer' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
309+
// expected-warning@+7{{'get_pointer<sycl::access::target::global_buffer, void>' is deprecated: accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
306310
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::global_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
307311
sycl::multi_ptr<int, sycl::access::address_space::global_space,
308312
sycl::access::decorated::legacy>
309313
LegacyGlobalMptr =
310314
sycl::make_ptr<int, sycl::access::address_space::global_space,
311315
sycl::access::decorated::legacy>(
312316
GlobalAcc.get_pointer());
317+
// expected-warning@+5{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
318+
// expected-warning@+7{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
319+
// expected-warning@+7{{'get_pointer' is deprecated: local_accessor::get_pointer() is deprecated, please use get_multi_ptr()}}
313320
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::local_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
314321
sycl::multi_ptr<int, sycl::access::address_space::local_space,
315322
sycl::access::decorated::legacy>
316323
LegacyLocalMptr =
317324
sycl::make_ptr<int, sycl::access::address_space::local_space,
318325
sycl::access::decorated::legacy>(
319326
LocalAcc.get_pointer());
320-
// expected-warning@+4{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
327+
328+
// expected-warning@+4{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
329+
// expected-warning@+5{{'make_ptr<int, sycl::access::address_space::private_space, sycl::access::decorated::legacy, std::enable_if<true>>' is deprecated: make_ptr is deprecated since SYCL 2020. Please use address_space_cast instead.}}
330+
// expected-warning@+6{{'legacy' is deprecated: sycl::access::decorated::legacy is deprecated since SYCL 2020}}
321331
sycl::multi_ptr<int, sycl::access::address_space::private_space,
322332
sycl::access::decorated::legacy>
323333
LegacyPrivateMptr =

0 commit comments

Comments
 (0)