[Openmp-commits] [PATCH] D148808: [OpenMP][libomptarget][AMDGPU] Enable optional active HSA wait state

[Ye Luo via Phabricator via Openmp-commits]

ye-luo added a comment.

Why separate KERNEL_BUSYWAIT and DATA_BUSYWAIT? I don't feel the need of two separate controls.
What are the units? Documentation?



================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:521
+        return Plugin::success();
+    }
     while (hsa_signal_wait_scacquire(Signal, HSA_SIGNAL_CONDITION_EQ, 0,
----------------
When reading the implementation of wait(). I'm still wondering how error got handled when kernels finished (signal equals 0) with an error.


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:1245
 
+  uint64_t KernelBusyWaitTics; // initialized from AMDGPUDeviceTy
+
----------------
Should this be private?


================
Comment at: openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp:1575
+  uint64_t KernelBusyWaitTics;
+  uint64_t DataBusyWaitTics;
+
----------------
Should they be private?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D148808