[OpenMP] Always print error messages in libomptarget CUDA plugin

[jhuber6]

Summary

Currently error messages from the CUDA plugins are only printed to the user if they have debugging enabled. Change this behaviour to always print the messages that result in offloading failure. This improves the error messages by indidcating what happened when the error occurs in the plugin library, such as a segmentation fault on the device.

[jdoerfert]

LGTM

openmp/libomptarget/plugins/cuda/src/rtl.cpp

 // Utility for retrieving and printing CUDA error string.
 #ifdef OMPTARGET_DEBUG
 #define CUDA_ERR_STRING(err)                                                   \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
       const char *errStr = nullptr;                                            \
       CUresult errStr_status = cuGetErrorString(err, &errStr);                 \
       if (errStr_status == CUDA_ERROR_INVALID_VALUE)                           \
-        DP("Unrecognized CUDA error code: %d\n", err);                         \
+        REPORT("Unrecognized CUDA error code: %d\n", err);                     \
       else if (errStr_status == CUDA_SUCCESS)                                  \
-        DP("CUDA error is: %s\n", errStr);                                     \
+        REPORT("CUDA error is: %s\n", errStr);                                 \
       else {                                                                   \
-        DP("Unresolved CUDA error code: %d\n", err);                           \
-        DP("Unsuccessful cuGetErrorString return status: %d\n",                \
-           errStr_status);                                                     \
+        REPORT("Unresolved CUDA error code: %d\n", err);                       \
+        REPORT("Unsuccessful cuGetErrorString return status: %d\n",            \
+               errStr_status);                                                 \
       }                                                                        \
+    } else {                                                                   \
+      const char *errStr = nullptr;                                            \
+      CUresult errStr_status = cuGetErrorString(err, &errStr);                 \
+      if (errStr_status == CUDA_SUCCESS)                                       \
+        REPORT("%s \n", errStr);                                               \
     }                                                                          \
   } while (false)
 #else // OMPTARGET_DEBUG
-#define CUDA_ERR_STRING(err) {}
+#define CUDA_ERR_STRING(err)                                                   \
+  do {                                                                         \
+    const char *errStr = nullptr;                                              \
+    CUresult errStr_status = cuGetErrorString(err, &errStr);                   \
+    if (errStr_status == CUDA_SUCCESS)                                         \
+      REPORT("%s \n", errStr);                                                 \
+  } while (false)
 #endif // OMPTARGET_DEBUG
 
 #include "../../common/elf_common.c"
 
 /// Keep entries table per device.
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
   std::vector<__tgt_offload_entry> Entries;
struct omptarget_device_environmentTy {
   int32_t debug_level;
 };
 
 namespace {
 bool checkResult(CUresult Err, const char *ErrMsg) {
   if (Err == CUDA_SUCCESS)
     return true;
 
-  DP("%s", ErrMsg);
+  REPORT("%s", ErrMsg);
   CUDA_ERR_STRING(Err);
   return false;
 }
 
 int memcpyDtoD(const void *SrcPtr, void *DstPtr, int64_t Size,
                CUstream Stream) {
   CUresult Err =
       cuMemcpyDtoDAsync((CUdeviceptr)DstPtr, (CUdeviceptr)SrcPtr, Size, Stream);
 
   if (Err != CUDA_SUCCESS) {
-    DP("Error when copying data from device to device. Pointers: src "
+    REPORT("Error when copying data from device to device. Pointers: src "
            "= " DPxMOD ", dst = " DPxMOD ", size = %" PRId64 "\n",
            DPxPTR(SrcPtr), DPxPTR(DstPtr), Size);
     CUDA_ERR_STRING(Err);
     return OFFLOAD_FAIL;
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 // Structure contains per-device data
int initDevice(const int DeviceId) {
 
     // Adjust teams to the env variables
     if (EnvTeamLimit > 0 && DeviceData[DeviceId].BlocksPerGrid > EnvTeamLimit) {
       DP("Capping max CUDA blocks per grid to OMP_TEAM_LIMIT=%d\n",
          EnvTeamLimit);
       DeviceData[DeviceId].BlocksPerGrid = EnvTeamLimit;
     }
 
-    if (getDebugLevel() || (getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))
-      INFO(DeviceId,
-           "Device supports up to %d CUDA blocks and %d threads with a "
-           "warp size of %d\n",
-           DeviceData[DeviceId].BlocksPerGrid,
-           DeviceData[DeviceId].ThreadsPerBlock, DeviceData[DeviceId].WarpSize);
 
     // Set default number of teams
     if (EnvNumTeams > 0) {
       DP("Default number of teams set according to environment %d\n",
          EnvNumTeams);
       DeviceData[DeviceId].NumTeams = EnvNumTeams;
     } else {
       DeviceData[DeviceId].NumTeams = DeviceRTLTy::DefaultNumTeams;
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
 
       if (E->size) {
         __tgt_offload_entry Entry = *E;
         CUdeviceptr CUPtr;
         size_t CUSize;
         Err = cuModuleGetGlobal(&CUPtr, &CUSize, Module, E->name);
         // We keep this style here because we need the name
         if (Err != CUDA_SUCCESS) {
-          DP("Loading global '%s' (Failed)\n", E->name);
+          REPORT("Loading global '%s' Failed\n", E->name);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
         if (CUSize != E->size) {
           DP("Loading global '%s' - size mismatch (%zd != %zd)\n", E->name,
              CUSize, E->size);
           return nullptr;
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
 
         continue;
       }
 
       CUfunction Func;
       Err = cuModuleGetFunction(&Func, Module, E->name);
       // We keep this style here because we need the name
       if (Err != CUDA_SUCCESS) {
-        DP("Loading '%s' (Failed)\n", E->name);
+        REPORT("Loading '%s' Failed\n", E->name);
         CUDA_ERR_STRING(Err);
         return nullptr;
       }
 
       DP("Entry point " DPxMOD " maps to %s (" DPxMOD ")\n",
          DPxPTR(E - HostBegin), E->name, DPxPTR(Func));
 
       // default value GENERIC (in case symbol is missing from cubin file)
for (const __tgt_offload_entry *E = HostBegin; E != HostEnd; ++E) {
         if (CUSize != sizeof(int8_t)) {
           DP("Loading global exec_mode '%s' - size mismatch (%zd != %zd)\n",
              ExecModeName, CUSize, sizeof(int8_t));
           return nullptr;
         }
 
         Err = cuMemcpyDtoH(&ExecModeVal, ExecModePtr, CUSize);
         if (Err != CUDA_SUCCESS) {
-          DP("Error when copying data from device to host. Pointers: "
+          REPORT("Error when copying data from device to host. Pointers: "
                  "host = " DPxMOD ", device = " DPxMOD ", size = %zd\n",
                  DPxPTR(&ExecModeVal), DPxPTR(ExecModePtr), CUSize);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
         if (ExecModeVal < 0 || ExecModeVal > 1) {
           DP("Error wrong exec_mode value specified in cubin file: %d\n",
              ExecModeVal);
           return nullptr;
         }
       } else {
-        DP("Loading global exec_mode '%s' - symbol missing, using default "
+        REPORT("Loading global exec_mode '%s' - symbol missing, using default "
                "value GENERIC (1)\n",
                ExecModeName);
         CUDA_ERR_STRING(Err);
       }
 
       KernelsList.emplace_back(Func, ExecModeVal);
 
       __tgt_offload_entry Entry = *E;
       Entry.addr = &KernelsList.back();
       addOffloadEntry(DeviceId, Entry);
#endif
 
       const char *DeviceEnvName = "omptarget_device_environment";
       CUdeviceptr DeviceEnvPtr;
       size_t CUSize;
 
       Err = cuModuleGetGlobal(&DeviceEnvPtr, &CUSize, Module, DeviceEnvName);
       if (Err == CUDA_SUCCESS) {
         if (CUSize != sizeof(DeviceEnv)) {
-          DP("Global device_environment '%s' - size mismatch (%zu != %zu)\n",
+          REPORT(
+              "Global device_environment '%s' - size mismatch (%zu != %zu)\n",
               DeviceEnvName, CUSize, sizeof(int32_t));
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
         Err = cuMemcpyHtoD(DeviceEnvPtr, &DeviceEnv, CUSize);
         if (Err != CUDA_SUCCESS) {
-          DP("Error when copying data from host to device. Pointers: "
+          REPORT("Error when copying data from host to device. Pointers: "
                  "host = " DPxMOD ", device = " DPxMOD ", size = %zu\n",
                  DPxPTR(&DeviceEnv), DPxPTR(DeviceEnvPtr), CUSize);
           CUDA_ERR_STRING(Err);
           return nullptr;
         }
 
         DP("Sending global device environment data %zu bytes\n", CUSize);
       } else {
         DP("Finding global device environment '%s' - symbol missing.\n",
            DeviceEnvName);
int dataSubmit(const int DeviceId, const void *TgtPtr, const void *HstPtr,
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
     if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
       return OFFLOAD_FAIL;
 
     CUstream Stream = getStream(DeviceId, AsyncInfoPtr);
 
     Err = cuMemcpyHtoDAsync((CUdeviceptr)TgtPtr, HstPtr, Size, Stream);
     if (Err != CUDA_SUCCESS) {
-      DP("Error when copying data from host to device. Pointers: host = " DPxMOD
-         ", device = " DPxMOD ", size = %" PRId64 "\n",
+      REPORT("Error when copying data from host to device. Pointers: host "
+             "= " DPxMOD ", device = " DPxMOD ", size = %" PRId64 "\n",
              DPxPTR(HstPtr), DPxPTR(TgtPtr), Size);
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }
 
     return OFFLOAD_SUCCESS;
   }
 
   int dataRetrieve(const int DeviceId, void *HstPtr, const void *TgtPtr,
                    const int64_t Size, __tgt_async_info *AsyncInfoPtr) const {
     assert(AsyncInfoPtr && "AsyncInfoPtr is nullptr");
 
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
     if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
       return OFFLOAD_FAIL;
 
     CUstream Stream = getStream(DeviceId, AsyncInfoPtr);
 
     Err = cuMemcpyDtoHAsync(HstPtr, (CUdeviceptr)TgtPtr, Size, Stream);
     if (Err != CUDA_SUCCESS) {
-      DP("Error when copying data from device to host. Pointers: host = " DPxMOD
-         ", device = " DPxMOD ", size = %" PRId64 "\n",
+      REPORT("Error when copying data from device to host. Pointers: host "
+             "= " DPxMOD ", device = " DPxMOD ", size = %" PRId64 "\n",
              DPxPTR(HstPtr), DPxPTR(TgtPtr), Size);
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }
 
     return OFFLOAD_SUCCESS;
   }
 
   int dataExchange(int SrcDevId, const void *SrcPtr, int DstDevId, void *DstPtr,
                    int64_t Size, __tgt_async_info *AsyncInfoPtr) const {
     assert(AsyncInfoPtr && "AsyncInfoPtr is nullptr");
 
     CUresult Err = cuCtxSetCurrent(DeviceData[SrcDevId].Context);
     if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
       return OFFLOAD_FAIL;
 
     CUstream Stream = getStream(SrcDevId, AsyncInfoPtr);
 
     // If they are two devices, we try peer to peer copy first
     if (SrcDevId != DstDevId) {
       int CanAccessPeer = 0;
       Err = cuDeviceCanAccessPeer(&CanAccessPeer, SrcDevId, DstDevId);
       if (Err != CUDA_SUCCESS) {
-        DP("Error returned from cuDeviceCanAccessPeer. src = %" PRId32
+        REPORT("Error returned from cuDeviceCanAccessPeer. src = %" PRId32
                ", dst = %" PRId32 "\n",
                SrcDevId, DstDevId);
         CUDA_ERR_STRING(Err);
         return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
       }
 
       if (!CanAccessPeer) {
         DP("P2P memcpy not supported so fall back to D2D memcpy");
         return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
       }
 
       Err = cuCtxEnablePeerAccess(DeviceData[DstDevId].Context, 0);
       if (Err != CUDA_SUCCESS) {
-        DP("Error returned from cuCtxEnablePeerAccess. src = %" PRId32
+        REPORT("Error returned from cuCtxEnablePeerAccess. src = %" PRId32
                ", dst = %" PRId32 "\n",
                SrcDevId, DstDevId);
         CUDA_ERR_STRING(Err);
         return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
       }
 
       Err = cuMemcpyPeerAsync((CUdeviceptr)DstPtr, DeviceData[DstDevId].Context,
                               (CUdeviceptr)SrcPtr, DeviceData[SrcDevId].Context,
                               Size, Stream);
       if (Err == CUDA_SUCCESS)
         return OFFLOAD_SUCCESS;
 
-      DP("Error returned from cuMemcpyPeerAsync. src_ptr = " DPxMOD
-         ", src_id =%" PRId32 ", dst_ptr = " DPxMOD ", dst_id =%" PRId32 "\n",
+      REPORT("Error returned from cuMemcpyPeerAsync. src_ptr = " DPxMOD
+             ", src_id =%" PRId32 ", dst_ptr = " DPxMOD ", dst_id =%" PRId32
+             "\n",
              DPxPTR(SrcPtr), SrcDevId, DPxPTR(DstPtr), DstDevId);
       CUDA_ERR_STRING(Err);
     }
 
     return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
   }
 
   int dataDelete(const int DeviceId, void *TgtPtr) const {
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
if (TeamNum <= 0) {
       DP("Capping number of teams to team limit %d\n",
          DeviceData[DeviceId].BlocksPerGrid);
       CudaBlocksPerGrid = DeviceData[DeviceId].BlocksPerGrid;
     } else {
       DP("Using requested number of teams %d\n", TeamNum);
       CudaBlocksPerGrid = TeamNum;
     }
 
-    if (getDebugLevel() || (getInfoLevel() & OMP_INFOTYPE_PLUGIN_KERNEL))
-      INFO(DeviceId,
-           "Launching kernel %s with %d blocks and %d threads in %s "
-           "mode\n",
-           (getOffloadEntry(DeviceId, TgtEntryPtr))
-               ? getOffloadEntry(DeviceId, TgtEntryPtr)->name
-               : "(null)",
-           CudaBlocksPerGrid, CudaThreadsPerBlock,
-           (KernelInfo->ExecutionMode == SPMD) ? "SPMD" : "Generic");
 
     CUstream Stream = getStream(DeviceId, AsyncInfo);
     Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
                          /* gridDimZ */ 1, CudaThreadsPerBlock,
                          /* blockDimY */ 1, /* blockDimZ */ 1,
                          /* sharedMemBytes */ 0, Stream, &Args[0], nullptr);
     if (!checkResult(Err, "Error returned from cuLaunchKernel\n"))
       return OFFLOAD_FAIL;
 
     DP("Launch of entry point at " DPxMOD " successful!\n",
        DPxPTR(TgtEntryPtr));
 
     return OFFLOAD_SUCCESS;
   }
 
   int synchronize(const int DeviceId, __tgt_async_info *AsyncInfoPtr) const {
     CUstream Stream = reinterpret_cast<CUstream>(AsyncInfoPtr->Queue);
     CUresult Err = cuStreamSynchronize(Stream);
     if (Err != CUDA_SUCCESS) {
-      DP("Error when synchronizing stream. stream = " DPxMOD
+      REPORT("Error when synchronizing stream. stream = " DPxMOD
              ", async info ptr = " DPxMOD "\n",
              DPxPTR(Stream), DPxPTR(AsyncInfoPtr));
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }
 
     // Once the stream is synchronized, return it to stream pool and reset
     // async_info. This is to make sure the synchronization only works for its
     // own tasks.
     StreamManager->returnStream(