[Openmp-commits] [PATCH] D146849: [OpenMP][libomptarget] Active and blocking HSA wait states

[Johannes Doerfert via Phabricator via Openmp-commits]

jdoerfert added inline comments.


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:501
+  AMDGPUSignalTy(uint64_t MicrosToWait)
+      : Signal({0}), UseCount(), MicrosToWait(MicrosToWait) {}
 
----------------
I would think we want a Env var and all Signal behave the same, so no argument here, wdyt?


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:558
+  /// Blocking the waiting thread
+  Error waitImpl() const {
+    // TODO: Is it better to use busy waiting or blocking the thread?
----------------



================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:570
+    hsa_signal_value_t Success = 0;
+    if (MicrosToWait) {
+      Got = hsa_signal_wait_scacquire(Signal, HSA_SIGNAL_CONDITION_NE, Init,
----------------
assert?


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:573-575
+      if (Got == Success) {
+        return Plugin::success();
+      }
----------------



Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D146849