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

[Kevin Sala via Phabricator via Openmp-commits]

kevinsala added inline comments.


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:568
+  Error activeWaitImpl(hsa_signal_value_t Init = 1) const {
+    hsa_signal_value_t Got = Init;
+    hsa_signal_value_t Success = 0;
----------------
Nit. This assignment is not needed, right? We could define it when it's used:

`hsa_signal_value_t Got = hsa_signal_wait_scacquire(...);`


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:571
+    if (MicrosToWait) {
+      Got = hsa_signal_wait_scacquire(Signal, HSA_SIGNAL_CONDITION_NE, Init,
+                                      MicrosToWait, HSA_WAIT_STATE_ACTIVE);
----------------
Could we make the exit condition consistent in both `activeWaitImpl` and `waitImpl`? The first is waiting until `Signal != Init (1)` and the second is waiting until `Signal == 0`.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D146849