[Openmp-commits] [openmp] [Libomptarget] Fix Nvidia offloading hanging on dataRetrieve using RPC (PR #66817)

[Joseph Huber via Openmp-commits]

================
@@ -0,0 +1,22 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: libc
+
+#include <assert.h>
+#include <stdio.h>
+
+#pragma omp declare target to(stdout)
+
+int main() {
+  int r = 0;
+// CHECK: PASS
+#pragma omp target map(from : r)
+  { r = fwrite("PASS\n", 1, sizeof("PASS\n") - 1, stdout); }
+  assert(r == sizeof("PASS\n") - 1 && "Incorrect number of bytes written");
----------------
jhuber6 wrote:

So, as far as I understand it, for a target region with a copy-back, we simply call this function to retrieve the data. Currently, NVPTX does an implicit sync because we schedule both the kernel and the copy-back on the steam stream. i.e. we cannot issue the memcpy until the kernel is done so we wait inside of the CUDA runtime. This is bad because it means that we aren't periodicaly checking the RPC server while waiting. 

https://github.com/llvm/llvm-project/pull/66817