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,7 +70,10 @@
 #define GETNAME2(name) #name
 #define GETNAME(name) GETNAME2(name)
 
-// Messaging interface
+// Generic interfaces for passing messages from the runtime to the user
+// Messages print notes about the implementation
+// Fatal and failure messages are used in the case of an error in the runtime
+// Info messages emit optional information messages about runtime values
 #define MESSAGE0(_str)                                                         \
   do {                                                                         \
     fprintf(stderr, GETNAME(TARGET_NAME) " message: %s\n", _str);              \
@@ -100,6 +103,12 @@
     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>
@@ -110,6 +119,7 @@
     fprintf(stderr, __VA_ARGS__);                                              \
   }
 
+// Emit a message for debugging
 #define DP(...)                                                                \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
@@ -117,6 +127,7 @@
     }                                                                          \
   } while (false)
 
+// Emit a message for debugging or failure if debugging is disabled
 #define REPORT(...)                                                            \
   do {                                                                         \
     if (getDebugLevel() > 0) {                                                 \
@@ -133,4 +144,12 @@
 #define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__);
 #endif // OMPTARGET_DEBUG
 
+// Emit a message giving the user extra information about the runtime
+#define INFO(...)                                                              \
+  do {                                                                         \
+    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);                                       \
@@ -495,6 +495,9 @@
     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) {
@@ -930,6 +933,15 @@
     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,
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;