[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

103

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

500

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?

943

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

936

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

936

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

936

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

298

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

82

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

openmp/libomptarget/include/Debug.h

#define FATAL_MESSAGE(_num, _str, ...)                                         \
   } while (0)
 
 #define FAILURE_MESSAGE(...)                                                   \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " error: ");                          \
     fprintf(stderr, __VA_ARGS__);                                              \
   } while (0)
 
+#define INFO_MESSAGE(...)                                                      \
+  do {                                                                         \
+    fprintf(stderr, GETNAME(TARGET_NAME) " info: ");                           \
+    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__);                                              \

 #else
 #define DEBUGP(prefix, ...)                                                    \
   {}
 #define DP(...)                                                                \
   {}
 #define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__);
 #endif // OMPTARGET_DEBUG
 
+#define INFO(...)                                                              \
+  do {                                                                         \
+    if (getInfoLevel() > 0) {                                                  \
+      INFO_MESSAGE(__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 {
   // 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();
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("Device max number of CUDA blocks %d, threads %d & warp size %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("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);
 
+    // Output information if enabled
+    if (KernelInfo->ExecutionMode == SPMD)
+      INFO("Launching kernel in SPMD mode\n");
+    else
+      INFO("Launching kernel in Generic mode\n");
+
+    INFO("Launching kernel with %d blocks and %d threads\n", CudaBlocksPerGrid,
+         CudaThreadsPerBlock);
+
     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;
 

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
+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;
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");
+          dumpTargetPointerMappings();
         FATAL_MESSAGE0(1, "failure of target construct while offloading is mandatory");
       }
       break;
   }
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
 EXTERN void __tgt_register_requires(int64_t flags) {
   RTLs->RegisterRequires(flags);
 }