Skip to content
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

Jump to bottom

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

Merged
merged 13 commits into from
Oct 14, 2022
Merged

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

merged 13 commits into from
Oct 14, 2022

Conversation

dongkyunahn-intel
Copy link
Contributor

  • __esimd_svm_atomic0/1/2
  • cmpxchng argument order fix (New value first, expected value second)
  • atomic_add/sub/min/max/cmpxchg update

@dongkyunahn-intel dongkyunahn-intel requested a review from a team as a code owner August 29, 2022 17:53
@dongkyunahn-intel
Copy link
Contributor Author

'dword_atomic_smoke.cpp' from intel/llvm-test-suite fails with infinite looping. Still debugging.

@dongkyunahn-intel dongkyunahn-intel mentioned this pull request Aug 29, 2022
Closed
@dongkyunahn-intel
Copy link
Contributor Author

'dword_atomic_smoke.cpp' from intel/llvm-test-suite fails with infinite looping. Still debugging.

The test is disabled in intel/llvm-test-suite as the test fails for GPU as well - intel/llvm-test-suite#1185

kbobrovs
kbobrovs reviewed Sep 2, 2022
View reviewed changes
sycl/include/sycl/ext/intel/esimd/detail/atomic_intrin.hpp Outdated
template <typename Ty> 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);
return __atomic_load_n((CmpxchgTy<Ty> *)ptr, __ATOMIC_SEQ_CST);
Copy link
Contributor

@kbobrovs kbobrovs Sep 2, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

type __atomic_load_n (type *ptr, int memorder)
so AFAIU return __atomic_load_n((CmpxchgTy<Ty> *)ptr, __ATOMIC_SEQ_CST); does type conversion from CmpxchgTy<Ty> to Ty before returning the result, which is wrong. Do you have tests for this? They should have caught this problem.

Suggested change
return __atomic_load_n((CmpxchgTy<Ty> *)ptr, __ATOMIC_SEQ_CST);
return sycl::bit_cast<Ty>(__atomic_load_n((CmpxchgTy<Ty> *)ptr, __ATOMIC_SEQ_CST));

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I don't have a unit test for this. I'll put TODO comments with __ESIMD_UNSUPPORTED_ON_HOST.

@dongkyunahn-intel
Copy link
Contributor Author

/verify with intel/llvm-test-suite#1259

@dongkyunahn-intel dongkyunahn-intel mentioned this pull request Sep 14, 2022
Merged
kbobrovs
kbobrovs reviewed Sep 16, 2022
View reviewed changes
sycl/include/sycl/ext/intel/experimental/esimd/detail/memory_intrin.hpp Outdated
@@ -1879,7 +1808,9 @@ __ESIMD_INTRIN void __esimd_lsc_fence(__ESIMD_DNS::simd_mask_storage_t<N> pred)
;
#else // __SYCL_DEVICE_ONLY__
{
__ESIMD_UNSUPPORTED_ON_HOST;
// In ESIMD_EMULATOR device interface, write operations are applied
Copy link
Contributor

Choose a reason for hiding this comment

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

This is extremely dangerous, and can be a source of obscure bugs.
Emulator uses multiple threads to emulate GPU threads, so memory fence can't be a no-op.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Fixed - built-in __atomic_thread_fence() is called.

@dongkyunahn-intel dongkyunahn-intel mentioned this pull request Sep 16, 2022
Merged
@dongkyunahn-intel
Copy link
Contributor Author

/verify with intel/llvm-test-suite#1274

1 similar comment
@dongkyunahn-intel
Copy link
Contributor Author

/verify with intel/llvm-test-suite#1274

@dongkyunahn-intel
Copy link
Contributor Author

Failures from ESIMD_EMULATOR are already handled in intel/llvm-test-suite#1274

********************
Timed Out Tests (1):
  SYCL :: dword_atomic_cmpxchg.cpp

********************
Failed Tests (3):
  SYCL :: lsc/lsc_slm.cpp
  SYCL :: lsc/lsc_surf.cpp
  SYCL :: lsc/lsc_usm.cpp

********************
Unexpectedly Passed Tests (1):
  SYCL :: dword_atomic_smoke.cpp

kbobrovs
kbobrovs reviewed Sep 19, 2022
View reviewed changes
sycl/include/sycl/ext/intel/esimd/detail/atomic_intrin.hpp
@@ -30,52 +39,87 @@ template <typename Ty> Ty atomic_store(Ty *ptr, Ty val) {
// 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

kbobrovs
kbobrovs reviewed Sep 19, 2022
View reviewed changes
sycl/include/sycl/ext/intel/esimd/detail/atomic_intrin.hpp
@@ -30,52 +39,87 @@ template <typename Ty> Ty atomic_store(Ty *ptr, Ty val) {
// 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);
__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.

kbobrovs
kbobrovs reviewed Sep 19, 2022
View reviewed changes
sycl/include/sycl/ext/intel/esimd/detail/atomic_intrin.hpp
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

@dongkyunahn-intel
Copy link
Contributor Author

lsc_slm/surf/usm will be fixed once argument order for cmpxchg is confirmed.

Other failures are handled in intel/llvm-test-suite#1274.

********************
Timed Out Tests (1):
  SYCL :: dword_atomic_cmpxchg.cpp

********************
Failed Tests (3):
  SYCL :: lsc/lsc_slm.cpp
  SYCL :: lsc/lsc_surf.cpp
  SYCL :: lsc/lsc_usm.cpp

********************
Unexpectedly Passed Tests (1):
  SYCL :: dword_atomic_smoke.cpp

@dongkyunahn-intel
Copy link
Contributor Author

/verify with intel/llvm-test-suite#1274

1 similar comment
@dongkyunahn-intel
Copy link
Contributor Author

/verify with intel/llvm-test-suite#1274

@dongkyunahn-intel
[SYCL][ESIMD][EMU] Atomic update
364874f 
- __esimd_svm_atomic0/1/2
- cmpxchng argument order fix (New value first, expected value second)
- atomic_add/sub/min/max/cmpxchg update
@dongkyunahn-intel
Typo fix
0197331
dongkyunahn-intel and others added 11 commits October 10, 2022 15:13
@dongkyunahn-intel
Removing duplicated declarations
b2cff06
@dongkyunahn-intel @kbobrovs
Apply suggestions from code review
a26feea 
Co-authored-by: kbobrovs <Konstantin.S.Bobrovsky@intel.com>
@dongkyunahn-intel
Addressing review comments
17e3d48 
- atomic intrinsic update and return value fix
@dongkyunahn-intel
Disable 'atomid_load' implementation
17f287a 
- To be enabled later with unit test
@dongkyunahn-intel
clang-format fix
61fab41
@dongkyunahn-intel
__esimd_lsc_fence fix
241bb93
@dongkyunahn-intel
Calling built-in __atomic_thread_fence() for *_fence() operations
7992db6
@dongkyunahn-intel
Work-around to avoid failure from odr.cpp test
4f33d98
@dongkyunahn-intel
Inlining atomic_fence instead of template argument for odr failure
fd9a2c3
@dongkyunahn-intel
Argument order fix for cmpxchg - comp value followed by new val
87e39d5
@dongkyunahn-intel
Adding 'inline' for 'atomic_*' functions
b6ae426 
- To prevent potential 'odr' failures from 'atomic_*' functions
instantiated with same template argument
@dongkyunahn-intel
Copy link
Contributor Author

Failures from 'SYCL / Linux / ESIMD Emu LLVM Test Suite (pull_request_target)'

********************
Timed Out Tests (1):
  SYCL :: dword_atomic_cmpxchg.cpp

********************
Unexpectedly Passed Tests (1):
  SYCL :: dword_atomic_smoke.cpp

These tests are modified in intel/llvm-test-suite#1274

@dongkyunahn-intel
Copy link
Contributor Author

/verify

@kbobrovs kbobrovs merged commit a6a0dea into intel:sycl Oct 14, 2022
@pvchupin pvchupin mentioned this pull request Oct 14, 2022
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@kbobrovs kbobrovs kbobrovs approved these changes

@v-klochkov v-klochkov Awaiting requested review from v-klochkov

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
2 participants
@dongkyunahn-intel @kbobrovs