diff --git a/openmp/libomptarget/include/Debug.h b/openmp/libomptarget/include/Debug.h
--- a/openmp/libomptarget/include/Debug.h
+++ b/openmp/libomptarget/include/Debug.h
@@ -70,23 +70,26 @@
 #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,   \
@@ -94,12 +97,20 @@
     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(...)                                                      \
+  do {                                                                         \
+    fprintf(stderr, GETNAME(TARGET_NAME) " info: ");                           \
+    fprintf(stderr, __VA_ARGS__);                                              \
+  } while (0)
+
 // Debugging messages
 #ifdef OMPTARGET_DEBUG
 #include <stdio.h>
@@ -110,6 +121,7 @@
     fprintf(stderr, __VA_ARGS__);                                              \
   }
 
+/// Emit a message for debugging if env LIBOMPTARGET_DEBUG > 0
 #define DP(...)                                                                \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
@@ -117,6 +129,7 @@
     }                                                                          \
   } while (false)
 
+/// Emit a message for debugging or error if debugging is disabled
 #define REPORT(...)                                                            \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
@@ -133,4 +146,15 @@
 #define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__);
 #endif // OMPTARGET_DEBUG
 
+/// Emit a message giving the user extra information about the runtime if 
+/// env LIBOMPTARGET_INFO > 0
+#define INFO(...)                                                              \
+  do {                                                                         \
+    if (getDebugLevel() > 0) {                                                 \
+      DEBUGP(DEBUG_PREFIX, __VA_ARGS__);                                       \
+    } else if (getInfoLevel() > 0) {                                           \
+      INFO_MESSAGE(__VA_ARGS__);                                               \
+    }                                                                          \
+  } while (false)
+
 #endif // _OMPTARGET_DEBUG_H
diff --git a/openmp/libomptarget/plugins/cuda/src/rtl.cpp b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
--- a/openmp/libomptarget/plugins/cuda/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/cuda/src/rtl.cpp
@@ -29,7 +29,7 @@
 #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);                                       \
@@ -492,9 +492,9 @@
       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 %d supports up to %d CUDA blocks and %d threads with a warp"
+         "size of %d\n", DeviceId, DeviceData[DeviceId].BlocksPerGrid,
+         DeviceData[DeviceId].ThreadsPerBlock, DeviceData[DeviceId].WarpSize);
 
     // Set default number of teams
     if (EnvNumTeams > 0) {
@@ -927,8 +927,9 @@
     }
 
     // Run on the device.
-    DP("Launch kernel with %d blocks and %d threads\n", CudaBlocksPerGrid,
-       CudaThreadsPerBlock);
+    INFO("Launching kernel with %d blocks and %d threads in %s mode\n",
+          CudaBlocksPerGrid, CudaThreadsPerBlock, 
+          (KernelInfo->ExecutionMode == SPMD)? "SPMD" : "Generic");
 
     CUstream Stream = getStream(DeviceId, AsyncInfo);
     Err = cuLaunchKernel(KernelInfo->Func, CudaBlocksPerGrid, /* gridDimY */ 1,
diff --git a/openmp/libomptarget/src/interface.cpp b/openmp/libomptarget/src/interface.cpp
--- a/openmp/libomptarget/src/interface.cpp
+++ b/openmp/libomptarget/src/interface.cpp
@@ -16,6 +16,7 @@
 #include "rtl.h"
 
 #include <cassert>
+#include <cstdio>
 #include <cstdlib>
 #include <mutex>
 
@@ -23,6 +24,20 @@
 kmp_target_offload_kind_t TargetOffloadPolicy = tgt_default;
 std::mutex TargetOffloadMtx;
 
+////////////////////////////////////////////////////////////////////////////////
+/// 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
 
@@ -61,7 +76,7 @@
     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;