Skip to content

Commit c84bbcd

Browse files
committed
fix
1 parent 1371754 commit c84bbcd

File tree

2 files changed

+24
-45
lines changed

2 files changed

+24
-45
lines changed

sycl/include/sycl/accessor.hpp

Lines changed: 23 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -363,9 +363,8 @@ class accessor_common {
363363
return MAccessor[MIDs];
364364
}
365365

366-
template <int CurDims = SubDims>
367-
typename std::enable_if_t<CurDims == 1 && IsAccessAtomic, atomic<DataT, AS>>
368-
operator[](size_t Index) const {
366+
template <int CurDims = SubDims,typename =std::enable_if_t<CurDims == 1 && IsAccessAtomic>>
367+
auto operator[](size_t Index) const {
369368
MIDs[Dims - CurDims] = Index;
370369
return MAccessor[MIDs];
371370
}
@@ -1727,39 +1726,34 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
17271726
return getQualifiedPtr()[LinearIndex];
17281727
}
17291728

1730-
template <int Dims = Dimensions>
1731-
operator typename std::enable_if_t<Dims == 0 &&
1732-
AccessMode == access::mode::atomic,
1733-
#ifdef __ENABLE_USM_ADDR_SPACE__
1734-
atomic<DataT>
1735-
#else
1736-
atomic<DataT, AS>
1737-
#endif
1738-
>() const {
1729+
template <int Dims = Dimensions,typename =std::enable_if_t<Dims == 0 && AccessMode == access::mode::atomic>>
1730+
operator
1731+
#ifdef __ENABLE_USM_ADDR_SPACE__
1732+
atomic<DataT>
1733+
#else
1734+
atomic<DataT, AS>
1735+
#endif
1736+
() const {
17391737
const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
1740-
return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
1738+
return Test(multi_ptr<DataT, AS, access::decorated::yes>(
17411739
getQualifiedPtr() + LinearIndex));
17421740
}
17431741

1744-
template <int Dims = Dimensions>
1745-
typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
1746-
atomic<DataT, AS>>
1747-
operator[](id<Dimensions> Index) const {
1742+
template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic>>
1743+
auto operator[](id<Dimensions> Index) const {
17481744
const size_t LinearIndex = getLinearIndex(Index);
17491745
return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
17501746
getQualifiedPtr() + LinearIndex));
17511747
}
17521748

1753-
template <int Dims = Dimensions>
1754-
typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
1755-
atomic<DataT, AS>>
1756-
operator[](size_t Index) const {
1749+
template <int Dims = Dimensions, typename = std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic>>
1750+
auto operator[](size_t Index) const {
17571751
const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
17581752
return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
17591753
getQualifiedPtr() + LinearIndex));
17601754
}
17611755
template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 1)>>
1762-
auto operator[](size_t Index) const {
1756+
AccessorSubscript<Dims - 1> operator[](size_t Index) const {
17631757
return AccessorSubscript<Dims - 1>(*this, Index);
17641758
}
17651759

@@ -2345,27 +2339,22 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
23452339
return getQualifiedPtr()[Index];
23462340
}
23472341

2348-
template <int Dims = Dimensions>
2349-
operator typename std::enable_if_t<
2350-
Dims == 0 && AccessMode == access::mode::atomic, atomic<DataT, AS>>()
2351-
const {
2342+
template <int Dims = Dimensions,typename=std::enable_if_t<
2343+
Dims == 0 && AccessMode == access::mode::atomic>>
2344+
operator atomic<DataT, AS>() const {
23522345
return atomic<DataT, AS>(
23532346
multi_ptr<DataT, AS, access::decorated::yes>(getQualifiedPtr()));
23542347
}
23552348

2356-
template <int Dims = Dimensions>
2357-
typename std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic,
2358-
atomic<DataT, AS>>
2359-
operator[](id<Dimensions> Index) const {
2349+
template <int Dims = Dimensions, typename = std::enable_if_t<(Dims > 0) && AccessMode == access::mode::atomic>>
2350+
auto operator[](id<Dimensions> Index) const {
23602351
const size_t LinearIndex = getLinearIndex(Index);
23612352
return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
23622353
getQualifiedPtr() + LinearIndex));
23632354
}
23642355

2365-
template <int Dims = Dimensions>
2366-
typename std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic,
2367-
atomic<DataT, AS>>
2368-
operator[](size_t Index) const {
2356+
template <int Dims = Dimensions, typename=std::enable_if_t<Dims == 1 && AccessMode == access::mode::atomic>>
2357+
auto operator[](size_t Index) const {
23692358
return atomic<DataT, AS>(multi_ptr<DataT, AS, access::decorated::yes>(
23702359
getQualifiedPtr() + Index));
23712360
}

sycl/include/sycl/atomic.hpp

Lines changed: 1 addition & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -166,11 +166,10 @@ extern T __spirv_AtomicMax(std::atomic<T> *Ptr, __spv::Scope::Flag,
166166

167167
namespace sycl {
168168
inline namespace _V1 {
169-
namespace detail {
170169

171170
template <typename T, access::address_space addressSpace =
172171
access::address_space::global_space>
173-
class atomic {
172+
class __SYCL2020_DEPRECATED("use sycl::atomic_ref instead") atomic {
174173
friend class atomic<T, access::address_space::global_space>;
175174
static_assert(detail::IsValidAtomicType<T>::value,
176175
"Invalid SYCL atomic type. Valid types are: int, "
@@ -400,15 +399,6 @@ T atomic_fetch_max(atomic<T, addressSpace> Object, T Operand,
400399
memory_order MemoryOrder = memory_order::relaxed) {
401400
return Object.fetch_max(Operand, MemoryOrder);
402401
}
403-
} // namespace detail
404-
// This alias is a workaround to allow accessor to use the deprecated
405-
// cl::sycl::atomic class without causing deprecation warnings without the user
406-
// explicitly using the corresponding accessor APIs.
407-
template <typename T, access::address_space addressSpace =
408-
access::address_space::global_space>
409-
using atomic __SYCL2020_DEPRECATED("use sycl::atomic_ref instead") =
410-
detail::atomic<T, addressSpace>;
411-
412402
} // namespace _V1
413403
} // namespace sycl
414404

0 commit comments

Comments
 (0)