[OpenMP] Begin Printing Information Dumps In Libomptarget and Plugins

[jhuber6]

Summary

This patch starts adding support for adding information dumps to libomptarget and rtl plugins. The information printing is controlled by the LIBOMPTARGET_INFO environment variable introduced in D86483. The goal of this patch is to provide the user with additional information about the device during kernel execution and providing the user with information dumps in the case of failure. This patch added the ability to dump the pointer mapping table as well as printing the number of blocks and threads in the cuda RTL

[ye-luo]

openmp/libomptarget/include/Debug.h

107

Please document all the macros, in this file, like functions. Say a bit about their intended use case.

openmp/libomptarget/src/interface.cpp

28

Update comments

[jhuber6]

Added additional comments. Should I add them to the doxygen notes at the top?

[ye-luo]

In D87165#2257164, @jhuber6 wrote:

Added additional comments. Should I add them to the doxygen notes at the top?

@jdoerfert Any suggestions?

I think the current documentation is insufficient. Adding/updating any documents for any of the macro or overall are helpful.
Although other macros are not touched, it is still helpful to document them like those you added. You are probably the best person who knows how and when such macros should be used now.

[jdoerfert]

In D87165#2257208, @ye-luo wrote:

In D87165#2257164, @jhuber6 wrote:

Added additional comments. Should I add them to the doxygen notes at the top?

@jdoerfert Any suggestions?

I think the current documentation is insufficient. Adding/updating any documents for any of the macro or overall are helpful.
Although other macros are not touched, it is still helpful to document them like those you added. You are probably the best person who knows how and when such macros should be used now.

I guess I'd prefer doxygen above the macro, as if they are functions. The reason is that I want us to make them functions at some point anyway.

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

499

Maybe: "Device %d supports up to % CUDA blocks and %d threads with a warp size of %d".
Also, why don't we make info behave like DP in debug mode and remove the DP macro?

948

Can we please merge those two messages:

INFO("Launching kernel with %d blocks and %d threads in %s mode\n", CudaBlocksPerGrid,
         CudaThreadsPerBlock,(KernelInfo->ExecutionMode == SPMD)? "SPMD" : "Generic")

openmp/libomptarget/src/interface.cpp

37

Don't we have to add space here to match the 18 above?

[jhuber6]

Updated messages and comments.

[jhuber6]

openmp/libomptarget/src/interface.cpp

37

The width of the pointer is 16 plus the "0x" at the start.

[jdoerfert]

I'm fine with this, one minor comment below. @ye-luo wdyt?

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

934

Forgot, we need to print the device number here too. Also the kernel name if we have it. That name is not great but has the line number in it (IIRC).

Maybe add a first mandatory argument to the INFO macro that is the device number. Then we can always forward something like: "[OMP-INFO][Device %d] " __VAR_ARGS__.
I see we already have the first part (to some degree) adding the device part would be good (I think).

[jhuber6]

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

934

You mean using the CUDA RTL to get the name from the CUfunction value in the KernelTy struct?

[jdoerfert]

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

934

We add it to KernelTy. When we do KernelsList.emplace_back(Func, ExecModeVal); in rtl.cpp, we know the name.

[jhuber6]

Added the target kernel name and device ID to the message. Added a new method to get the offload entry given the pointer. I can't decide if it's better to just always print the device ID if we have it versus putting it in the macro itself.

Does this require any test files?

[jdoerfert]

In D87165#2261026, @jhuber6 wrote:

Added the target kernel name and device ID to the message. Added a new method to get the offload entry given the pointer. I can't decide if it's better to just always print the device ID if we have it versus putting it in the macro itself.

I believe there is no reason to be somewhere in this methods without a device ID. Thus, I would require it as part of the function/macro.

Does this require any test files?

I think that would be good. We should be able to set the INFO variable in the run line as well. Do we have a test for the DP macro we can copy?

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

299

can we replace findOffloadEntry with this please.

[ye-luo]

The changes I requested as been added. Remove my blocking. Still need other reviews to be addressed.

[jhuber6]

Added small test file.

[jhuber6]

Changed table dumps to require LIBOMPTARGET_INFO > 1. We could move this to a bitvector in the future. Added an extra message to inform users to set LIBOMPTARGET_INFO on failure.

[jdoerfert]

LGTM

[ye-luo]

openmp/libomptarget/src/interface.cpp

81

Need to be more specific about the "tables"
host-device memory mapping tables?

openmp/libomptarget/include/Debug.h

 #include <inttypes.h>
 #undef __STDC_FORMAT_MACROS
 
 #define DPxMOD "0x%0*" PRIxPTR
 #define DPxPTR(ptr) ((int)(2 * sizeof(uintptr_t))), ((uintptr_t)(ptr))
 #define GETNAME2(name) #name
 #define GETNAME(name) GETNAME2(name)
 
-// Messaging interface
+/// Print a generic message string from libomptarget or a plugin RTL
 #define MESSAGE0(_str)                                                         \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " message: %s\n", _str);              \
   } while (0)
 
+/// Print a printf formatting string message from libomptarget or a plugin RTL
 #define MESSAGE(_str, ...)                                                     \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " message: " _str "\n", __VA_ARGS__); \
   } while (0)
 
+/// Print fatal error message with an error string and error identifier
 #define FATAL_MESSAGE0(_num, _str)                                             \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " fatal error %d: %s\n", _num, _str); \
     abort();                                                                   \
   } while (0)
 
+/// Print fatal error message with a printf string and error identifier
 #define FATAL_MESSAGE(_num, _str, ...)                                         \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " fatal error %d:" _str "\n", _num,   \
             __VA_ARGS__);                                                      \
     abort();                                                                   \
   } while (0)
 
+/// Print a generic error string from libomptarget or a plugin RTL
 #define FAILURE_MESSAGE(...)                                                   \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " error: ");                          \
     fprintf(stderr, __VA_ARGS__);                                              \
   } while (0)
 
+/// Print a generic information string used if LIBOMPTARGET_INFO=1
+#define INFO_MESSAGE(_num, ...)                                                \
+  do {                                                                         \
+    fprintf(stderr, GETNAME(TARGET_NAME) " device %d info: ", _num);           \
+    fprintf(stderr, __VA_ARGS__);                                              \
+  } while (0)
+
 // Debugging messages
 #ifdef OMPTARGET_DEBUG
 #include <stdio.h>
 
 #define DEBUGP(prefix, ...)                                                    \
   {                                                                            \
     fprintf(stderr, "%s --> ", prefix);                                        \
     fprintf(stderr, __VA_ARGS__);                                              \
   }
 
+/// Emit a message for debugging
 #define DP(...)                                                                \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
       DEBUGP(DEBUG_PREFIX, __VA_ARGS__);                                       \
     }                                                                          \
   } while (false)
 
+/// Emit a message for debugging or failure if debugging is disabled
 #define REPORT(...)                                                            \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
       DP(__VA_ARGS__);                                                         \
     } else {                                                                   \
       FAILURE_MESSAGE(__VA_ARGS__);                                            \
     }                                                                          \
   } while (false)
 #else
 #define DEBUGP(prefix, ...)                                                    \
   {}
 #define DP(...)                                                                \
   {}
 #define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__);
 #endif // OMPTARGET_DEBUG
 
+/// Emit a message giving the user extra information about the runtime if
+#define INFO(_id, ...)                                                         \
+  do {                                                                         \
+    if (getDebugLevel() > 0) {                                                 \
+      DEBUGP(DEBUG_PREFIX, __VA_ARGS__);                                       \
+    } else if (getInfoLevel() > 0) {                                           \
+      INFO_MESSAGE(_id, __VA_ARGS__);                                          \
+    }                                                                          \
+  } while (false)
+
 #endif // _OMPTARGET_DEBUG_H

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

 
 #define TARGET_NAME CUDA
 #define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
 
 // Utility for retrieving and printing CUDA error string.
 #ifdef OMPTARGET_DEBUG
 #define CUDA_ERR_STRING(err)                                                   \
   do {                                                                         \
-    if (getDebugLevel() > 0) {                                                      \
+    if (getDebugLevel() > 0) {                                                 \
       const char *errStr;                                                      \
       cuGetErrorString(err, &errStr);                                          \
       DP("CUDA error is: %s\n", errStr);                                       \
     }                                                                          \
   } while (false)
 #else // OMPTARGET_DEBUG
 #define CUDA_ERR_STRING(err) {}
 #endif // OMPTARGET_DEBUG
class DeviceRTLTy {
   std::vector<CUmodule> Modules;
 
   // Record entry point associated with device
   void addOffloadEntry(const int DeviceId, const __tgt_offload_entry entry) {
     FuncOrGblEntryTy &E = DeviceData[DeviceId].FuncGblEntries.back();
     E.Entries.push_back(entry);
   }
 
-  // Return true if the entry is associated with device
-  bool findOffloadEntry(const int DeviceId, const void *Addr) const {
+  // Return a pointer to the entry associated with the pointer
+  const __tgt_offload_entry *getOffloadEntry(const int DeviceId,
+                                             const void *Addr) const {
     for (const __tgt_offload_entry &Itr :
          DeviceData[DeviceId].FuncGblEntries.back().Entries)
       if (Itr.addr == Addr)
-        return true;
+        return &Itr;
 
-    return false;
+    return nullptr;
   }
 
   // Return the pointer to the target entries table
   __tgt_target_table *getOffloadEntriesTable(const int DeviceId) {
     FuncOrGblEntryTy &E = DeviceData[DeviceId].FuncGblEntries.back();
 
     if (E.Entries.empty())
       return nullptr;
 
     // Update table info according to the entries and return the pointer
     E.Table.EntriesBegin = E.Entries.data();
     E.Table.EntriesEnd = E.Entries.data() + E.Entries.size();
 
     return &E.Table;
   }
 
   // Clear entries table for a device
   void clearOffloadEntriesTable(const int DeviceId) {
     DeviceData[DeviceId].FuncGblEntries.emplace_back();
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;
     }
 
-    DP("Max number of CUDA blocks %d, threads %d & warp size %d\n",
-       DeviceData[DeviceId].BlocksPerGrid, DeviceData[DeviceId].ThreadsPerBlock,
-       DeviceData[DeviceId].WarpSize);
+    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;
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;
     }
 
-    // Run on the device.
-    DP("Launch kernel with %d blocks and %d threads\n", CudaBlocksPerGrid,
-       CudaThreadsPerBlock);
+    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);

openmp/libomptarget/src/interface.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "device.h"
 #include "private.h"
 #include "rtl.h"
 
 #include <cassert>
+#include <cstdio>
 #include <cstdlib>
 #include <mutex>
 
 // Store target policy (disabled, mandatory, default)
 kmp_target_offload_kind_t TargetOffloadPolicy = tgt_default;
 std::mutex TargetOffloadMtx;
 
 ////////////////////////////////////////////////////////////////////////////////
-/// manage the success or failure of a target construct
+/// dump a table of all the host-target pointer pairs on failure
+static void dumpTargetPointerMappings() {
+  for (const auto &Device : Devices) {
+    fprintf(stderr, "Device %d:\n", Device.DeviceID);
+    fprintf(stderr, "%-18s %-18s %s\n", "Host Ptr", "Target Ptr", "Size (B)");
+    for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
+      fprintf(stderr, DPxMOD " " DPxMOD " %lu\n",
+              DPxPTR(HostTargetMap.HstPtrBegin),
+              DPxPTR(HostTargetMap.TgtPtrBegin),
+              HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin);
+    }
+  }
+}
 
+////////////////////////////////////////////////////////////////////////////////
+/// manage the success or failure of a target construct
 static void HandleDefaultTargetOffload() {
   TargetOffloadMtx.lock();
   if (TargetOffloadPolicy == tgt_default) {
     if (omp_get_num_devices() > 0) {
       DP("Default TARGET OFFLOAD policy is now mandatory "
          "(devices were found)\n");
       TargetOffloadPolicy = tgt_mandatory;
     } else {
case tgt_disabled:
       }
       break;
     case tgt_default:
       FATAL_MESSAGE0(1, "default offloading policy must be switched to "
                         "mandatory or disabled");
       break;
     case tgt_mandatory:
       if (!success) {
-        if (getInfoLevel() > 0)
-          MESSAGE0("LIBOMPTARGET_INFO is not supported yet");
+        if (getInfoLevel() > 1)
+          dumpTargetPointerMappings();
+        else
+          FAILURE_MESSAGE("run with env LIBOMPTARGET_INFO>1 to dump tables\n");
+
         FATAL_MESSAGE0(1, "failure of target construct while offloading is mandatory");
       }
       break;
   }
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags

openmp/libomptarget/test/offloading/info.c

+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_INFO=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
+
+#include <stdio.h>
+#include <omp.h>
+
+int main() {
+    int ptr = 1;
+
+// INFO: CUDA device {{[0-9]+}} info: Device supports up to {{[0-9]+}} CUDA blocks and {{[0-9]+}} threads with a warp size of {{[0-9]+}}
+// INFO: CUDA device {{[0-9]+}} info: Launching kernel {{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
+#pragma omp target map(tofrom:ptr)
+  {ptr = 1;}
+
+  return 0;
+}