[Openmp-commits] [PATCH] D154172: [OpenMP] Added memory scope to atomic::inc API and used the device scope in reduction.

[Johannes Doerfert via Phabricator via Openmp-commits]

jdoerfert added inline comments.


================
Comment at: openmp/libomptarget/DeviceRTL/include/Synchronization.h:33
+  cgroup  // All threads in the contention group, e.g. the team
+};
+
----------------
Can we out a namespace or sth around this, `atomic::device` and `atomic::all` are confusing.
What about `atomic::scope::ABC`?


================
Comment at: openmp/libomptarget/DeviceRTL/src/Synchronization.cpp:207
+  }
+}
+
----------------
The above is expanded below.


================
Comment at: openmp/libomptarget/DeviceRTL/src/Synchronization.cpp:213-226
   switch (Ordering) {
   default:
     __builtin_unreachable();
   case atomic::relaxed:
-    return __builtin_amdgcn_atomic_inc32(A, V, atomic::relaxed, "");
+    return atomicIncRelaxed(A, V, MemScope);
   case atomic::aquire:
+    return atomicIncAquire(A, V, MemScope);
----------------





Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D154172