[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:571
+    if (MicrosToWait) {
+      Got = hsa_signal_wait_scacquire(Signal, HSA_SIGNAL_CONDITION_NE, Init,
+                                      MicrosToWait, HSA_WAIT_STATE_ACTIVE);
----------------
jplehr wrote:
> 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?
That new condition looks good to me. But we should extend the code documentation of these two functions to specify the new behavior. The wait operation does no longer stop when the signal is observed with a zero value, but when we stop observing the `Init` (default: 1) value.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D146849