[Openmp-commits] [PATCH] D145831: [OpenMP][libomptarget] Add support for critical regions in AMD GPU device offloading

[Gheorghe-Teodor Bercea via Phabricator via Openmp-commits]

doru1004 added inline comments.
Herald added a subscriber: jplehr.


================
Comment at: openmp/libomptarget/DeviceRTL/src/Synchronization.cpp:287
+    while (!atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::relaxed,
+                      atomic::relaxed) != UNSET) {
+      __builtin_amdgcn_s_sleep(32);
----------------
jdoerfert wrote:
> Shouldn't the release and acquire semantics be part of this CAS?
> If we run relaxed who is to say we see the update of another thread. I would have assumed on failure we want an acquire fence and on success a release fence. Not sure if we need additional ones.
The idea is that the thread only executes the acquire once when it enters the critical region to prevent it from being executed at every iteration of the loop (which comes with performance penalties if it happens).

I would like to understand why you say there's a chance we don't see the update of another thread. The unsetting of the lock is happening here:
```
void unsetLock(omp_lock_t *Lock) {
  (void)atomicExchange((uint32_t *)Lock, UNSET, atomic::acq_rel);
}
```


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D145831/new/

https://reviews.llvm.org/D145831