[OpenMP][libomptarget] Fixed an issue that device sync is skipped if the kernel doesn't have any argument

[tianshilei1992]

Summary

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.

[jdoerfert]

LGTM, one NIT below.

openmp/libomptarget/src/omptarget.cpp

1323

", and it is invoked asynchronously", really? Maybe the language is just confusing.

openmp/libomptarget/src/omptarget.cpp

struct DeallocTgtPtrInfo {
   /// Whether it has \p close modifier
   bool HasCloseModifier;
 
   DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool ForceDelete,
                     bool HasCloseModifier)
       : 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.
 int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
                   void **ArgBases, void **Args, int64_t *ArgSizes,
                   int64_t *ArgTypes, map_var_info_t *ArgNames,
                   void **ArgMappers, __tgt_async_info *AsyncInfo) {
   int Ret;
int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
 
   // We need to synchronize before deallocating data.
   // If AsyncInfo is nullptr, the previous data transfer (if has) will be
   // synchronous, so we don't need to synchronize again. If AsyncInfo->Queue is
   // nullptr, there is no data transfer happened because once there is,
   // 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
   for (DeallocTgtPtrInfo &Info : DeallocTgtPtrs) {
     Ret = Device.deallocTgtPtr(Info.HstPtrBegin, Info.DataSize,
                                Info.ForceDelete, Info.HasCloseModifier);
     if (Ret != OFFLOAD_SUCCESS) {
       REPORT("Deallocating data from device failed.\n");
       return OFFLOAD_FAIL;
if (ArgNum) {
     // variables
     Ret = processDataAfter(loc, DeviceId, HostPtr, ArgNum, ArgBases, Args,
                            ArgSizes, ArgTypes, ArgNames, ArgMappers,
                            PrivateArgumentManager, &AsyncInfo);
     if (Ret != OFFLOAD_SUCCESS) {
       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 the device supports async operations, so we need a
+    // sync at this point.
+    return syncDevice(Device, &AsyncInfo);
   }
 
   return OFFLOAD_SUCCESS;
 }

openmp/libomptarget/test/offloading/assert.cpp

+// 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;
+}