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

[Jan-Patrick Lehr via Phabricator via Openmp-commits]

jplehr marked an inline comment as done.
jplehr added inline comments.


================
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);
----------------
kevinsala wrote:
> 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`.
Sure.
I would vote for the `!=` version. though I'd need to see what changes are necessary. Basically, making the `wait` implementation logically the same as in AOMP: https://github.com/RadeonOpenCompute/llvm-project/blob/amd-stg-open/openmp/libomptarget/plugins/amdgpu/impl/impl.cpp#L16

Or are there objections?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D146849