[Openmp-commits] [PATCH] D96067: [OpenMP][libomptarget] Fixed an issue that device sync is skipped if the kernel doesn't have any argument

[Shilei Tian via Phabricator via Openmp-commits]

tianshilei1992 created this revision.
tianshilei1992 added reviewers: jdoerfert, JonChesterfield.
Herald added subscribers: guansong, yaxunl.
tianshilei1992 requested review of this revision.
Herald added subscribers: openmp-commits, sstefan1.
Herald added a project: OpenMP.

Currently if there is not kernel argument, device synchronization will
be skipped. This can lead to two issues:

1. If there is any device error, it will not be captured;
2. The target region might end before the kernel is done, which is not spec conformant.

The test added in this patch only runs on NVPTX platform, although it will not
be executed by Phab at all. It also requires `not` which is not available on most
systems.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D96067

Files:
  openmp/libomptarget/src/omptarget.cpp
  openmp/libomptarget/test/offloading/assert.cpp


Index: openmp/libomptarget/test/offloading/assert.cpp
===================================================================
--- /dev/null
+++ openmp/libomptarget/test/offloading/assert.cpp
@@ -0,0 +1,8 @@
+// RUN: %libomptarget-compilexx-nvptx64-nvidia-cuda && %libomptarget-run-fail-nvptx64-nvidia-cuda
+
+int main(int argc, char *argv[]) {
+#pragma omp target
+  { __builtin_trap(); }
+
+  return 0;
+}
Index: openmp/libomptarget/src/omptarget.cpp
===================================================================
--- openmp/libomptarget/src/omptarget.cpp
+++ openmp/libomptarget/src/omptarget.cpp
@@ -451,6 +451,17 @@
       : HstPtrBegin(HstPtr), DataSize(Size), ForceDelete(ForceDelete),
         HasCloseModifier(HasCloseModifier) {}
 };
+
+/// Synchronize device
+static int syncDevice(DeviceTy &Device, __tgt_async_info *AsyncInfo) {
+  assert(AsyncInfo && AsyncInfo->Queue && "Invalid AsyncInfo");
+  if (Device.synchronize(AsyncInfo) != OFFLOAD_SUCCESS) {
+    REPORT("Failed to synchronize device.\n");
+    return OFFLOAD_FAIL;
+  }
+
+  return OFFLOAD_SUCCESS;
+}
 } // namespace
 
 /// Internal function to undo the mapping and retrieve the data from the device.
@@ -631,11 +642,9 @@
   // AsyncInfo->Queue will not be nullptr, so again, we don't need to
   // synchronize.
   if (AsyncInfo && AsyncInfo->Queue) {
-    Ret = Device.synchronize(AsyncInfo);
-    if (Ret != OFFLOAD_SUCCESS) {
-      REPORT("Failed to synchronize device.\n");
+    Ret = syncDevice(Device, AsyncInfo);
+    if (Ret != OFFLOAD_SUCCESS)
       return OFFLOAD_FAIL;
-    }
   }
 
   // Deallocate target pointer
@@ -1307,6 +1316,11 @@
       REPORT("Failed to process data after launching the kernel.\n");
       return OFFLOAD_FAIL;
     }
+  } else if (AsyncInfo.Queue) {
+    // If ArgNum is zero, but AsyncInfo.Queue is valid, then the kernel doesn't
+    // hava any argument, and it is invoked asynchronously. We need a sync at
+    // this point.
+    return syncDevice(Device, &AsyncInfo);
   }
 
   return OFFLOAD_SUCCESS;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D96067.321525.patch
Type: text/x-patch
Size: 2024 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20210204/6d5060d6/attachment.bin>