[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 marked an inline comment as done.
doru1004 added inline comments.


================
Comment at: openmp/libomptarget/DeviceRTL/src/Synchronization.cpp:285
+    while (!atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::relaxed,
+                      atomic::relaxed) != UNSET) {
+      __builtin_amdgcn_s_sleep(32);
----------------
dhruvachak wrote:
> dhruvachak wrote:
> > I think the CAS should take an acq_rel as the success memorder and an acq as the failure memorder. 
> > 
> > The acq in both the cases will ensure it sees an update from another thread. The release from the other updating thread is not sufficient for this thread to see that update. And when this thread succeeds, the release will ensure other threads see the update of this thread. 
> > 
> > With the above memorders, I would think we could get rid of the fences.
> > I think the CAS should take an acq_rel as the success memorder and an acq as the failure memorder. 
> > 
> > The acq in both the cases will ensure it sees an update from another thread. The release from the other updating thread is not sufficient for this thread to see that update. And when this thread succeeds, the release will ensure other threads see the update of this thread. 
> > 
> > With the above memorders, I would think we could get rid of the fences.
> 
> We had an offline discussion. Based on that, I am ok with the current memorders in the patch for the atomicCAS in setLock. C++11 spec has this statement "Implementations should make atomic stores visible to atomic loads within a reasonable amount of time." and since the CAS is done in a loop, the setLock should eventually see the update made by another thread that executed the unsetLock.
Sounds good. I'll mark as resolved.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D145831