Skip to content

Commit

Permalink
SYCL: Replace deprecated atomic types and operations (AMReX-Codes#2921)
Browse files Browse the repository at this point in the history 
* SYCL: Replace deprecated atomic types and operations

* Change atomic refs to device memory scope

When using the relaxed memory order, the memory scope is ignored.
Thus, for cosmetic reasons only, we set the memory scope to device, the broadest option when using the global address space.

Co-authored-by: Weiqun Zhang <WeiqunZhang@lbl.gov>
  • Loading branch information
@nmnobre @WeiqunZhang
nmnobre and WeiqunZhang authored Sep 4, 2022
1 parent cc3cd14 commit fb0b31e
Showing 1 changed file with 37 additions and 26 deletions.
63 changes: 37 additions & 26 deletions Src/Base/AMReX_GpuAtomic.H
View file
Original file line number Diff line number Diff line change
Expand Up @@ -30,15 +30,16 @@ namespace detail {
{
#if defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
static_assert(sizeof(R) == sizeof(I), "sizeof R != sizeof I");
I* const add_as_I = reinterpret_cast<I*>(address);
sycl::atomic<I,as> a{sycl::multi_ptr<I,as>(add_as_I)};
I old_I = a.load(mo), new_I;
sycl::atomic_ref<I,mo,ms,as> a{*add_as_I};
I old_I = a.load(), new_I;
do {
R const new_R = f(*(reinterpret_cast<R const*>(&old_I)), val);
new_I = *(reinterpret_cast<I const*>(&new_R));
} while (! a.compare_exchange_strong(old_I, new_I, mo));
} while (! a.compare_exchange_strong(old_I, new_I));
return *(reinterpret_cast<R const*>(&old_I));
#else
R old = *address;
Expand All @@ -53,17 +54,18 @@ namespace detail {
{
#if defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
static_assert(sizeof(R) == sizeof(I), "sizeof R != sizeof I");
I* const add_as_I = reinterpret_cast<I*>(address);
sycl::atomic<I, as> a{sycl::multi_ptr<I,as>(add_as_I)};
I old_I = a.load(mo), new_I;
sycl::atomic_ref<I,mo,ms,as> a{*add_as_I};
I old_I = a.load(), new_I;
bool test_success;
do {
R const tmp = op(*(reinterpret_cast<R const*>(&old_I)), val);
new_I = *(reinterpret_cast<I const*>(&tmp));
test_success = cond(tmp);
} while (test_success && ! a.compare_exchange_strong(old_I, new_I, mo));
} while (test_success && ! a.compare_exchange_strong(old_I, new_I));
return test_success;
#else
R old = *address;
Expand Down Expand Up @@ -131,9 +133,10 @@ namespace detail {
return atomicAdd(sum, value);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(sum)};
return a.fetch_add(value, mo);
sycl::atomic_ref<T,mo,ms,as> a{*sum};
return a.fetch_add(value);
#else
amrex::ignore_unused(sum, value);
return T(); // should never get here, but have to return something
Expand Down Expand Up @@ -313,9 +316,10 @@ namespace detail {
return atomicMin(m, value);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(m)};
return a.fetch_min(value, mo);
sycl::atomic_ref<T,mo,ms,as> a{*m};
return a.fetch_min(value);
#else
amrex::ignore_unused(m,value);
return T(); // should never get here, but have to return something
Expand Down Expand Up @@ -373,9 +377,10 @@ namespace detail {
return atomicMax(m, value);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(m)};
return a.fetch_max(value, mo);
sycl::atomic_ref<T,mo,ms,as> a{*m};
return a.fetch_max(value);
#else
amrex::ignore_unused(m,value);
return T(); // should never get here, but have to return something
Expand Down Expand Up @@ -430,9 +435,10 @@ namespace detail {
return atomicOr(m, value);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<int,as> a{sycl::multi_ptr<int,as>(m)};
return a.fetch_or(value, mo);
sycl::atomic_ref<int,mo,ms,as> a{*m};
return a.fetch_or(value);
#else
int const old = *m;
*m = (*m) || value;
Expand All @@ -451,9 +457,10 @@ namespace detail {
return atomicAnd(m, value ? ~0x0 : 0);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<int,as> a{sycl::multi_ptr<int,as>(m)};
return a.fetch_and(value ? ~0x0 : 0, mo);
sycl::atomic_ref<int,mo,ms,as> a{*m};
return a.fetch_and(value ? ~0x0 : 0);
#else
int const old = *m;
*m = (*m) && value;
Expand All @@ -472,11 +479,12 @@ namespace detail {
{
#if defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
sycl::atomic<unsigned int,AS> a{sycl::multi_ptr<unsigned int,AS>(m)};
unsigned int oldi = a.load(mo), newi;
constexpr auto ms = sycl::memory_scope::device;
sycl::atomic_ref<unsigned int,mo,ms,AS> a{*m};
unsigned int oldi = a.load(), newi;
do {
newi = (oldi >= value) ? 0u : (oldi+1u);
} while (! a.compare_exchange_strong(oldi, newi, mo));
} while (! a.compare_exchange_strong(oldi, newi));
return oldi;
#else
auto const old = *m;
Expand Down Expand Up @@ -509,12 +517,13 @@ namespace detail {
return atomicDec(m, value);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<unsigned int,as> a{sycl::multi_ptr<unsigned int,as>(m)};
unsigned int oldi = a.load(mo), newi;
sycl::atomic_ref<unsigned int,mo,ms,as> a{*m};
unsigned int oldi = a.load(), newi;
do {
newi = ((oldi == 0u) || (oldi > value)) ? value : (oldi-1u);
} while (! a.compare_exchange_strong(oldi, newi, mo));
} while (! a.compare_exchange_strong(oldi, newi));
return oldi;
#else
auto const old = *m;
Expand All @@ -535,9 +544,10 @@ namespace detail {
return atomicExch(address, val);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(address)};
return sycl::atomic_exchange(a, val, mo);
sycl::atomic_ref<T,mo,ms,as> a{*address};
return a.exchange(val);
#else
auto const old = *address;
*address = val;
Expand All @@ -557,9 +567,10 @@ namespace detail {
return atomicCAS(address, compare, val);
#elif defined(__SYCL_DEVICE_ONLY__)
constexpr auto mo = sycl::memory_order::relaxed;
constexpr auto ms = sycl::memory_scope::device;
constexpr auto as = sycl::access::address_space::global_space;
sycl::atomic<T,as> a{sycl::multi_ptr<T,as>(address)};
a.compare_exchange_strong(compare, val, mo);
sycl::atomic_ref<T,mo,ms,as> a{*address};
a.compare_exchange_strong(compare, val);
return compare;
#else
auto const old = *address;
Expand Down

0 comments on commit fb0b31e

Please sign in to comment.