Skip to content

[SYCL][ESIMD][EMU] Atomic update #6661

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 13 commits into from
Oct 14, 2022
Merged

[SYCL][ESIMD][EMU] Atomic update #6661

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
13 commits
Select commit Hold shift + click to select a range
364874f
[SYCL][ESIMD][EMU] Atomic update
dongkyunahn-intel Aug 26, 2022
0197331
Typo fix
dongkyunahn-intel Aug 29, 2022
b2cff06
Removing duplicated declarations
dongkyunahn-intel Aug 29, 2022
a26feea
Apply suggestions from code review
dongkyunahn-intel Aug 30, 2022
17e3d48
Addressing review comments
dongkyunahn-intel Aug 30, 2022
17f287a
Disable 'atomid_load' implementation
dongkyunahn-intel Sep 6, 2022
61fab41
clang-format fix
dongkyunahn-intel Sep 13, 2022
241bb93
__esimd_lsc_fence fix
dongkyunahn-intel Sep 13, 2022
7992db6
Calling built-in __atomic_thread_fence() for *_fence() operations
dongkyunahn-intel Sep 16, 2022
4f33d98
Work-around to avoid failure from odr.cpp test
dongkyunahn-intel Sep 16, 2022
fd9a2c3
Inlining atomic_fence instead of template argument for odr failure
dongkyunahn-intel Sep 19, 2022
87e39d5
Argument order fix for cmpxchg - comp value followed by new val
dongkyunahn-intel Sep 22, 2022
b6ae426
Adding 'inline' for 'atomic_*' functions
dongkyunahn-intel Sep 23, 2022
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
135 changes: 107 additions & 28 deletions sycl/include/sycl/ext/intel/esimd/detail/atomic_intrin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,126 +16,205 @@ namespace __ESIMD_DNS {
// This function implements atomic update of pre-existing variable in the
// absense of C++ 20's atomic_ref.

template <typename Ty> Ty atomic_load(Ty *ptr) {
// __atomic_* functions support only integral types. In order to
// support floating types for certain operations like min/max,
// 'cmpxchg' operation is applied for result values using
// 'bridging' variables in integral type.
template <typename Ty> using CmpxchgTy = __ESIMD_DNS::uint_type_t<sizeof(Ty)>;

template <typename Ty> inline Ty atomic_load(Ty *ptr) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
return __atomic_load(ptr, __ATOMIC_SEQ_CST);
__ESIMD_UNSUPPORTED_ON_HOST;
// TODO : Enable with unit test
/* return sycl::bit_cast<Ty>(__atomic_load_n((CmpxchgTy<Ty> *)ptr,
__ATOMIC_SEQ_CST)); */
#endif
}

template <typename Ty> Ty atomic_store(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_store(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
__atomic_store(ptr, val, __ATOMIC_SEQ_CST);
Ty ret = atomic_load<Ty>((CmpxchgTy<Ty> *)ptr);
Copy link
Contributor

Choose a reason for hiding this comment

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

it looks strange that we need atomic_load as a part of atomic_store implementation, but that's how Gen ISA is :)
no action for this comment is needed

__atomic_store_n((CmpxchgTy<Ty> *)ptr, val, __ATOMIC_SEQ_CST);
Copy link
Contributor

Choose a reason for hiding this comment

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

Nit: strictly speaking, reinterpret_cast should be used. Can be fixed later.

return ret;
#endif
}

template <typename Ty> Ty atomic_add_fetch(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_add(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
return __atomic_add_fetch(ptr, val, __ATOMIC_SEQ_CST);
if constexpr (std::is_integral_v<Ty>) {
return __atomic_fetch_add(ptr, val, __ATOMIC_SEQ_CST);
} else {
// For Floating type
Ty _old, _new;
CmpxchgTy<Ty> _old_bits, _new_bits;
do {
_old = *ptr;
_new = _old + val;
_old_bits = *(CmpxchgTy<Ty> *)&_old;
_new_bits = *(CmpxchgTy<Ty> *)&_new;
} while (!__atomic_compare_exchange_n((CmpxchgTy<Ty> *)ptr, &_old_bits,
_new_bits, false, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST));
return _old;
}
#endif
}

template <typename Ty> Ty atomic_sub_fetch(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_sub(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
return __atomic_sub_fetch(ptr, val, __ATOMIC_SEQ_CST);
if constexpr (std::is_integral_v<Ty>) {
return __atomic_fetch_sub(ptr, val, __ATOMIC_SEQ_CST);
} else {
// For Floating type
Ty _old, _new;
CmpxchgTy<Ty> _old_bits, _new_bits;
do {
_old = *ptr;
_new = _old - val;
_old_bits = *(CmpxchgTy<Ty> *)&_old;
_new_bits = *(CmpxchgTy<Ty> *)&_new;
} while (!__atomic_compare_exchange_n((CmpxchgTy<Ty> *)ptr, &_old_bits,
_new_bits, false, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST));
return _old;
}
#endif
}

template <typename Ty> Ty atomic_and_fetch(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_and(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
return __atomic_and_fetch(ptr, val, __ATOMIC_SEQ_CST);
static_assert(std::is_integral<Ty>::value);
return __atomic_fetch_and(ptr, val, __ATOMIC_SEQ_CST);
#endif
}

template <typename Ty> Ty atomic_or_fetch(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_or(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
return __atomic_or_fetch(ptr, val, __ATOMIC_SEQ_CST);
static_assert(std::is_integral<Ty>::value);
return __atomic_fetch_or(ptr, val, __ATOMIC_SEQ_CST);
#endif
}

template <typename Ty> Ty atomic_xor_fetch(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_xor(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
return __atomic_xor_fetch(ptr, val, __ATOMIC_SEQ_CST);
static_assert(std::is_integral<Ty>::value);
return __atomic_fetch_xor(ptr, val, __ATOMIC_SEQ_CST);
#endif
}

template <typename Ty> Ty atomic_min(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_min(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
// TODO FIXME: fix implementation for FP types.
if constexpr (std::is_integral_v<Ty>) {
Ty _old, _new;
do {
_old = *ptr;
_new = std::min<Ty>(_old, val);
} while (!__atomic_compare_exchange_n(ptr, &_old, _new, false,
__ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST));
return _new;
return _old;
} else {
__ESIMD_UNSUPPORTED_ON_HOST;
Ty _old, _new;
CmpxchgTy<Ty> _old_bits, _new_bits;
do {
_old = *ptr;
_new = std::min(_old, val);
_old_bits = *(CmpxchgTy<Ty> *)&_old;
_new_bits = *(CmpxchgTy<Ty> *)&_new;
} while (!__atomic_compare_exchange_n((CmpxchgTy<Ty> *)ptr, &_old_bits,
_new_bits, false, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST));
return _old;
}
#endif
}

template <typename Ty> Ty atomic_max(Ty *ptr, Ty val) {
template <typename Ty> inline Ty atomic_max(Ty *ptr, Ty val) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
// TODO FIXME: fix implementation for FP types.
if constexpr (std::is_integral_v<Ty>) {
Ty _old, _new;
do {
_old = *ptr;
_new = std::max<Ty>(_old, val);
} while (!__atomic_compare_exchange_n(ptr, &_old, _new, false,
__ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST));
return _new;
return _old;
} else {
__ESIMD_UNSUPPORTED_ON_HOST;
Ty _old, _new;
CmpxchgTy<Ty> _old_bits, _new_bits;
do {
_old = *ptr;
_new = std::max(_old, val);
_old_bits = *(CmpxchgTy<Ty> *)&_old;
Copy link
Contributor

@kbobrovs kbobrovs Sep 19, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

Here and in many more places:
Nit: sycl::bitcast should be used instead. Can be fixed later

_new_bits = *(CmpxchgTy<Ty> *)&_new;
} while (!__atomic_compare_exchange_n((CmpxchgTy<Ty> *)(CmpxchgTy<Ty> *)ptr,
&_old_bits, _new_bits, false,
__ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST));
return _old;
}
#endif
}

template <typename Ty> Ty atomic_cmpxchg(Ty *ptr, Ty expected, Ty desired) {
template <typename Ty>
inline Ty atomic_cmpxchg(Ty *ptr, Ty expected, Ty desired) {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
// TODO FIXME: fix implementation for FP types.
if constexpr (std::is_integral_v<Ty>) {
Ty _old = expected;
__atomic_compare_exchange_n(ptr, &_old, desired, false, __ATOMIC_SEQ_CST,
Ty local = expected;
__atomic_compare_exchange_n(ptr, &local, desired, false, __ATOMIC_SEQ_CST,
__ATOMIC_SEQ_CST);
return *ptr;
// if exchange occured, this means 'local=expected=*ptr'. So local
// is returned as old val
// if exchange did not occur, *ptr value compared against 'local'
// is stored in 'local'. So local is returned as old val
return local;
} else {
__ESIMD_UNSUPPORTED_ON_HOST;
CmpxchgTy<Ty> desired_bits = *(CmpxchgTy<Ty> *)&desired;
CmpxchgTy<Ty> local_bits = *(CmpxchgTy<Ty> *)&expected;
__atomic_compare_exchange_n((CmpxchgTy<Ty> *)ptr, &local_bits, desired_bits,
false, __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST);
return *((Ty *)&local_bits);
}
#endif
}

inline void atomic_fence() {
#ifdef _WIN32
// TODO: Windows will be supported soon
__ESIMD_UNSUPPORTED_ON_HOST;
#else
__atomic_thread_fence(__ATOMIC_SEQ_CST);
#endif
}

} // namespace __ESIMD_DNS

/// @endcond ESIMD_DETAIL
91 changes: 75 additions & 16 deletions sycl/include/sycl/ext/intel/esimd/detail/memory_intrin.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -485,7 +485,25 @@ __esimd_svm_atomic0(__ESIMD_DNS::vector_type_t<uint64_t, N> addrs,
;
#else
{
__ESIMD_UNSUPPORTED_ON_HOST;
__ESIMD_DNS::vector_type_t<Ty, N> Oldval = 0;

for (int AddrIdx = 0; AddrIdx < N; AddrIdx += 1) {
if (pred[AddrIdx] == 0) {
// Skip Oldval vector elements correpsonding to
// predicates whose value is zero
continue;
}
if constexpr (Op == __ESIMD_NS::atomic_op::load) {
Oldval[AddrIdx] = __ESIMD_DNS::atomic_load<Ty>((Ty *)addrs[AddrIdx]);
} else if constexpr (Op == __ESIMD_NS::atomic_op::inc) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_add<Ty>((Ty *)addrs[AddrIdx], static_cast<Ty>(1));
} else if constexpr (Op == __ESIMD_NS::atomic_op::dec) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_sub<Ty>((Ty *)addrs[AddrIdx], static_cast<Ty>(1));
}
}
return Oldval;
}
#endif // __SYCL_DEVICE_ONLY__

Expand All @@ -498,23 +516,49 @@ __esimd_svm_atomic1(__ESIMD_DNS::vector_type_t<uint64_t, N> addrs,
;
#else
{
__ESIMD_DNS::vector_type_t<Ty, N> retv;
__ESIMD_DNS::vector_type_t<Ty, N> Oldval;

for (int i = 0; i < N; i++) {
if (pred[i]) {
Ty *p = reinterpret_cast<Ty *>(addrs[i]);
for (int AddrIdx = 0; AddrIdx < N; AddrIdx++) {
if (pred[AddrIdx] == 0) {
// Skip Output vector elements correpsonding to
// predicates whose value is zero
continue;
}

switch (Op) {
case __ESIMD_NS::atomic_op::add:
retv[i] = __ESIMD_DNS::atomic_add_fetch<Ty>(p, src0[i]);
break;
default:
__ESIMD_UNSUPPORTED_ON_HOST;
}
if constexpr (Op == __ESIMD_NS::atomic_op::store) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_store<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr ((Op == __ESIMD_NS::atomic_op::add) ||
(Op == __ESIMD_NS::atomic_op::fadd)) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_add<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr ((Op == __ESIMD_NS::atomic_op::sub) ||
(Op == __ESIMD_NS::atomic_op::fsub)) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_sub<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr ((Op == __ESIMD_NS::atomic_op::minsint) ||
(Op == __ESIMD_NS::atomic_op::min) ||
(Op == __ESIMD_NS::atomic_op::fmin)) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_min<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr ((Op == __ESIMD_NS::atomic_op::maxsint) ||
(Op == __ESIMD_NS::atomic_op::max) ||
(Op == __ESIMD_NS::atomic_op::fmax)) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_max<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr (Op == __ESIMD_NS::atomic_op::bit_and) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_and<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr (Op == __ESIMD_NS::atomic_op::bit_or) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_or<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
} else if constexpr (Op == __ESIMD_NS::atomic_op::bit_xor) {
Oldval[AddrIdx] =
__ESIMD_DNS::atomic_xor<Ty>((Ty *)addrs[AddrIdx], src0[AddrIdx]);
}
}

return retv;
return Oldval;
}
#endif // __SYCL_DEVICE_ONLY__

Expand All @@ -528,7 +572,20 @@ __esimd_svm_atomic2(__ESIMD_DNS::vector_type_t<uint64_t, N> addrs,
;
#else
{
__ESIMD_UNSUPPORTED_ON_HOST;
__ESIMD_DNS::vector_type_t<Ty, N> Oldval;

for (int AddrIdx = 0; AddrIdx < N; AddrIdx++) {
if (pred[AddrIdx] == 0) {
// Skip Output vector elements correpsonding to
// predicates whose value is zero
continue;
}
static_assert((Op == __ESIMD_NS::atomic_op::cmpxchg) ||
(Op == __ESIMD_NS::atomic_op::fcmpxchg));
Oldval[AddrIdx] = __ESIMD_DNS::atomic_cmpxchg((Ty *)addrs[AddrIdx],
src0[AddrIdx], src1[AddrIdx]);
}
return Oldval;
}
#endif // __SYCL_DEVICE_ONLY__

Expand Down Expand Up @@ -557,7 +614,9 @@ __ESIMD_INTRIN void __esimd_fence(uint8_t cntl)
;
#else
{
sycl::detail::getESIMDDeviceInterface()->cm_fence_ptr();
// CM_EMU's 'cm_fence' is NOP. Disabled.
// sycl::detail::getESIMDDeviceInterface()->cm_fence_ptr();
__ESIMD_DNS::atomic_fence();
}
#endif // __SYCL_DEVICE_ONLY__

Expand Down Expand Up @@ -849,7 +908,7 @@ __esimd_dword_atomic0(__ESIMD_DNS::simd_mask_storage_t<N> pred,

switch (Op) {
case __ESIMD_NS::atomic_op::inc:
retv[i] = __ESIMD_DNS::atomic_add_fetch<Ty>(p, 1);
retv[i] = __ESIMD_DNS::atomic_add<Ty>(p, 1);
break;
default:
__ESIMD_UNSUPPORTED_ON_HOST;
Expand Down
Loading