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

[Dhruva Chakrabarti via Phabricator via Openmp-commits]

dhruvachak added inline comments.


================
Comment at: openmp/libomptarget/DeviceRTL/src/Synchronization.cpp:283
+  if (mapping::getThreadIdInWarp() == LowestActiveThread) {
+    fenceKernel(atomic::release);
+    while (!atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::relaxed,
----------------
Why do we need the release fence in setLock before the CAS?


================
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:
> 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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D145831