[OpenMP] Add using bit flags to select Libomptarget Information

[jhuber6]

Summary

This patch adds more fine-grained support over which information is output from the libomptarget runtime when run with the environment variable LIBOMPTARGET_INFO set. An extensible set of flags can be used to pick and choose which information the user is interested in.

[JonChesterfield]

I like the direction. Looks easily applicable to amdgpu and granular debug printing is a good feature.

As commented above, not convinced the -1 sentinel is worth the complexity. If the goal is to only call getenv once (which seems sensible), that's probably worth guarding with something threadsafe instead.

openmp/libomptarget/include/Debug.h

41

Why not uint32, with zero for all flags clear as the default?

[jhuber6]

openmp/libomptarget/include/Debug.h

41

I wanted to have a method for getting the info / debug levels that didn't expose them as global variables, so I put them in a function as static variables. But this required having a special value for uninitialized so the value is only read from getenv the first time it's called. I could change it to use a boolean or something and avoid the complexity of reserving the highest order bit.

[JonChesterfield]

openmp/libomptarget/include/Debug.h

41

How about code that runs on library load? We have a few getenv calls and it's probably good to do all of them once.

struct env {
env() {all the getenv calls}

uint32_t debug_level() {return dl;}
uint32_t info_level() {}

private:
  uint32_t dl; // etc
} environment;

Makes the variables read-only after construction. Or use const member variables.

That will be thread safe, removes some branches, no sentinel. No teardown problems as destructor is a no-op.

[jhuber6]

Changing the info and debug functions to use booleans to check if they've been initialized already. Another benefit is that allows every flag to be enabled by setting LIBOMPTARGET_INFO=1.

[jhuber6]

openmp/libomptarget/include/Debug.h

41

This is basically how it was done previously, I didn't like needing to initialize it in every plugin separately when they're all sharing the same function. I could change it to work this way but I feel like the new approach using an extra boolean works alright.

[JonChesterfield]

openmp/libomptarget/include/Debug.h

41

That is worth avoiding. The variable in the header doesn't achieve that though - each plugin will have it's own copy of the static variable + wrapper function.

If we move the init code, however it is done, into libomptarget then each plugin will need to make a cross-dl call to get the variable, but getenv is only called once. I think that's a win.

A comprehensively overengineered solution is to make the getenv calls once, from libomptarget, and also have the plugin read&cache the result of said call.

[jdoerfert]

@JonChesterfield Let's do this firt and we can revisit a single global in libomptarget later. WDYT?

[JonChesterfield]

Sure, may as well. If there's a race here it's preexisting.

[JonChesterfield]

Oh - didn't mention it's still signed. Would prefer move to uint32 as the -1 is no longer in use, but that's not blocking.

[jhuber6]

Changing info and debug values to uint32_t.

[JonChesterfield]

Nice, thanks!

[grokos]

openmp/libomptarget/include/Debug.h

63

Is 0x0008 reserved for something else?

68

I would suggest you declare initialized as std::once_flag and call std::call_once() to initialize InfoLevel. This is guaranteed to be thread-safe (unlike the implementation posted here) while also making sure that InfoLevel is initialized exactly once. We initialize RTLs in RTLsTy::RegisterLib() in the very same way.

openmp/libomptarget/src/interface.cpp

237

Wouldn't we want to have

if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
  printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
                       arg_names);

here as well? It is useful to see what arguments we have if something goes wrong at a target exit data directive or upon exiting a target data scope.

[jhuber6]

openmp/libomptarget/include/Debug.h

63

No real reason, I just felt like keeping a space open, so if we add another libomptarget flag the first nibble just applies to libomptarget, since the fourth one is for the CUDA plugin right now.

68

Seems like a good solution. Should I push this first and address that in a future patch? Or just do it here.

openmp/libomptarget/src/interface.cpp

237

This method just prints the arguments to __tgt_target_data_begin_mapper and __tgt_target_data_end_mapper, which should have the same arguments right? I had it in previously but it didn't add any useful information when I tested it. I could add a separate message that just says something like "Exiting data region started at (loc)". Would that be good?

[jhuber6]

Changing environment variables to initialize using std::call_once. Added a message when exiting an OpenMP data kernel.

[grokos]

openmp/libomptarget/src/interface.cpp

237

For scoped constructs (e.g. omp target data) this is true, whatever arguments we have in begin_mapper the very same arguments we have in end_mapper. But what if we have a standalone directive (e.g. omp target exit data)? If I recall correctly (currently on vacation with no access to a working llvm environment), omp target exit data is compiled into a __tgt_target_data_end_mapper call. In this case there may be no matching enter data directive so if we don't call printKernelArguments() here the user will not be able to see what arguments end_mapper was called with.

[jhuber6]

openmp/libomptarget/src/interface.cpp

237

That's true, I should probably just make it print whether its a data region entry, device kernel, or data region exit. That way it's clear what the flow is. Otherwise it would be confusing because it would still say that a variable is mapped "to" at the end of a data region.

[jhuber6]

Fixing a bug with passing a nullptr and making the kernel arguments print a different message for which type of region was executed.

[grokos]

Thanks for the changes! LGTM.

[vzakhari]

openmp/libomptarget/include/Debug.h

72–81

Hello. Am I the only one who is getting libomptarget.rtl.x86_64.so build fail with -DCMAKE_BUILD_TYPE=Debug?

ld: CMakeFiles/omptarget.rtl.x86_64.dir/__/generic-elf-64bit/src/rtl.cpp.o: in function `getInfoLevel()::{lambda()#1}::operator()() const':
.../openmp/libomptarget/include/Debug.h:61: undefined reference to `getInfoLevel()::InfoLevel'
ld: CMakeFiles/omptarget.rtl.x86_64.dir/__/generic-elf-64bit/src/rtl.cpp.o: relocation R_X86_64_PC32 against undefined symbol `_ZZL12getInfoLevelvE9InfoLevel' can not be used when making a shared object; recompile with -fPIC

Here is a small reproducer:

#include <mutex>

static inline uint32_t getInfoLevel() {
  static uint32_t InfoLevel = 0;
  static std::once_flag Flag{};
  std::call_once(Flag, []() {
      InfoLevel = 1;
  });

  return InfoLevel;
}

I cannot build a shared library from it using gcc: g++ -fPIC -std=c++14 test.cpp -shared -o test.so

ld: /tmp/ccBpTXXs.o: relocation R_X86_64_PC32 against undefined symbol `_ZZL12getInfoLevelvE9InfoLevel' can not be used when making a shared object; recompile with -fPIC

The issue is reproducible with GCC 5.5.0 and 6.5.0, and does not reproduce with 7.3.0. I believe LLVM's minimal requirement is 5.1, so it is a problem.

The issue seems to be caused by the fact that getInfoLevel() is not referenced in openmp/libomptarget/plugins/generic-elf-64bit/src/rtl.cpp

[jhuber6]

openmp/libomptarget/include/Debug.h

72–81

I have no issues compiling the reproducer you provided. It compiles without errors locally and here https://godbolt.org/z/oEKq3E for the compilers you listed. There might be a problem when the compiler doesn't create a static variable when the function in unused. Does the problem go away if you create a dummy usage of getInfoLevel()?

[jhuber6]

openmp/libomptarget/include/Debug.h

72–81

Actually you're right, it doesn't work on Godbolt, the error went off screen and didn't list as an error in the compilation. It definitely works locally on the machine I tested it on so who knows what's up with that. The error goes away if I just add a dummy usage to force the compiler to actually generate a symbol for InfoLevel. It's a solution but probably not the best.

[vzakhari]

openmp/libomptarget/include/Debug.h

72–81

Of course, it does compile with a dummy reference. I guess you have more recent GCC on your machine. The issue seems to be fixed somewhere between 6.5.0 and 7.1.0.

[grokos]

I tried inserting the following forward-declaration right before the function definition and the problem seems to go away:

static inline uint32_t getInfoLevel() __attribute__((used));

I guess it's better to have the above rather than a dummy call?

[JonChesterfield]

Can mark the function itself as ((used)), for slightly less typing. Seems a reasonable workaround to me, with a comment like 'workaround bug in gcc 5.ish'

[jhuber6]

In D93727#2523719, @JonChesterfield wrote:

Can mark the function itself as ((used)), for slightly less typing. Seems a reasonable workaround to me, with a comment like 'workaround bug in gcc 5.ish'

Yeah, that's an option. Adding a dummy variable would defeat the original purpose of initializing the variables this way, rather than having each plugin initialize it individually and globally. This problem only seems to occur on GCC x86 from the compilers I've checked. The reason It wasn't failing on my machine was because it was GCC 6.3 on PPC. I'm assuming any other compilers will just ignore the attribute if they don't understand it. This should be sufficient.

static inline uint32_t __attribute__((used)) getInfoLevel  () { ... }
static inline uint32_t __attribute__((used)) getDebugLevel  () { ... }

[vzakhari]

In D93727#2523723, @jhuber6 wrote:

In D93727#2523719, @JonChesterfield wrote:

Can mark the function itself as ((used)), for slightly less typing. Seems a reasonable workaround to me, with a comment like 'workaround bug in gcc 5.ish'

Yeah, that's an option. Adding a dummy variable would defeat the original purpose of initializing the variables this way, rather than having each plugin initialize it individually and globally. This problem only seems to occur on GCC x86 from the compilers I've checked. The reason It wasn't failing on my machine was because it was GCC 6.3 on PPC. I'm assuming any other compilers will just ignore the attribute if they don't understand it. This should be sufficient.

static inline uint32_t __attribute__((used)) getInfoLevel  () { ... }
static inline uint32_t __attribute__((used)) getDebugLevel  () { ... }

Thank you all for the investigation. Fixed in https://reviews.llvm.org/D95486

[vzakhari]

offloading/dynamic_module_load.c fails due to this commit in my environment with x86_64 offload. It fails because std::call_once cannot be used without libpthread, and libomptarget.rtl.x86_64.so is not linked with libpthread. Basically, gdb shows that include/c++/7.3.0/x86_64-linux-gnu/bits/gthr-default.h::__gthread_active_p() returns false, when std::call_once is called from libomptarget.rtl.x86_64.so. It does return true, though, when std::call_once is called from libomptarget.so, because libomptarget.so is linked with libpthread.

It looks like GCC11 will have a fix: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=55394
For the time being, all plugins have to be linked with libpthread, since Debug.h is considered to be common code for all plugins.

Any disagreement?

[jdoerfert]

In D93727#2524396, @vzakhari wrote:

offloading/dynamic_module_load.c fails due to this commit in my environment with x86_64 offload. It fails because std::call_once cannot be used without libpthread, and libomptarget.rtl.x86_64.so is not linked with libpthread. Basically, gdb shows that include/c++/7.3.0/x86_64-linux-gnu/bits/gthr-default.h::__gthread_active_p() returns false, when std::call_once is called from libomptarget.rtl.x86_64.so. It does return true, though, when std::call_once is called from libomptarget.so, because libomptarget.so is linked with libpthread.

It looks like GCC11 will have a fix: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=55394
For the time being, all plugins have to be linked with libpthread, since Debug.h is considered to be common code for all plugins.

Any disagreement?

Linking libpthread also into the plugins seems plausible. Can you write the patch?

openmp/libomptarget/include/Debug.h

 //                               // 8 digits for 32bit systems
 //                               // 16 digits for 64bit
 //   (uintptr_t) ptr);
 //
 //===----------------------------------------------------------------------===//
 #ifndef _OMPTARGET_DEBUG_H
 #define _OMPTARGET_DEBUG_H
 
-static inline int getInfoLevel() {
-  static int InfoLevel = -1;
+/// 31-Bit field data attributes controlling information presented to the user.
+enum OpenMPInfoType : int32_t {
+  // Print data arguments and attributes upon entering an OpenMP device kernel.
+  OMP_INFOTYPE_KERNEL_ARGS = 0x0001,
+  // Indicate when an address already exists in the device mapping table.
+  OMP_INFOTYPE_MAPPING_EXISTS = 0x0002,
+  // Dump the contents of the device pointer map at kernel exit or failure.
+  OMP_INFOTYPE_DUMP_TABLE = 0x0004,
+  // Print kernel information from target device plugins
+  OMP_INFOTYPE_PLUGIN_KERNEL = 0x0010,
+};
+
+static inline int32_t getInfoLevel() {
+  static int32_t InfoLevel = -1;
   if (InfoLevel >= 0)
     return InfoLevel;
 
   if (char *EnvStr = getenv("LIBOMPTARGET_INFO"))
     InfoLevel = std::stoi(EnvStr);
+  else
+    InfoLevel = 0;
 
   return InfoLevel;
 }
 
-static inline int getDebugLevel() {
-  static int DebugLevel = -1;
+static inline int32_t getDebugLevel() {
+  static int32_t DebugLevel = -1;
   if (DebugLevel >= 0)
     return DebugLevel;
 
   if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))
     DebugLevel = std::stoi(EnvStr);
+  else
+    DebugLevel = 0;
 
   return DebugLevel;
 }
 
 #ifndef __STDC_FORMAT_MACROS
 #define __STDC_FORMAT_MACROS
 #endif
 #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)
 
 /// Print a generic message 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, GETNAME(TARGET_NAME) " device %d info: ", (int)_num);      \
     fprintf(stderr, __VA_ARGS__);                                              \
   } while (0)
 
 // Debugging messages
 #ifdef OMPTARGET_DEBUG
 #include <stdio.h>
 
 #define DEBUGP(prefix, ...)                                                    \

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

 //===----RTLs/cuda/src/rtl.cpp - Target RTLs Implementation ------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // RTL for CUDA machine
 //
 //===----------------------------------------------------------------------===//
 
 #include <cassert>
 #include <cstddef>
 #include <cuda.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <string>
 #include <vector>
 
 #include "Debug.h"
 #include "omptargetplugin.h"
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;
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;

openmp/libomptarget/src/device.cpp

: DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false), InitFlag(),
       HasPendingGlobals(false), HostDataToTargetMap(), PendingCtorsDtors(),
       ShadowPtrMap(), DataMapMtx(), PendingGlobalsMtx(), ShadowMtx(),
       MemoryManager(nullptr) {}
 
 DeviceTy::~DeviceTy() {
   if (DeviceID == -1 || getInfoLevel() < 1)
     return;
 
-  dumpTargetPointerMappings(*this);
+  dumpTargetPointerMappings(nullptr, *this);
 }
 
 int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
   DataMapMtx.lock();
 
   // Check if entry exists
   auto search = HostDataToTargetMap.find(HstPtrBeginTy{(uintptr_t)HstPtrBegin});
   if (search != HostDataToTargetMap.end()) {
if (lr.Flags.IsContained ||
       ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && IsImplicit)) {
     auto &HT = *lr.Entry;
     IsNew = false;
 
     if (UpdateRefCount)
       HT.incRefCount();
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+    if (getDebugLevel() || getInfoLevel() & OMP_INFOTYPE_MAPPING_EXISTS)
       INFO(DeviceID,
            "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
            ", "
            "Size=%" PRId64 ",%s RefCount=%s, Name=%s\n",
            (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(tp),
            Size, (UpdateRefCount ? " updated" : ""),
-         HT.isRefCountInf() ? "INF" : std::to_string(HT.getRefCount()).c_str(),
+           HT.isRefCountInf() ? "INF"
+                              : std::to_string(HT.getRefCount()).c_str(),
            (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "(null)");
     rc = (void *)tp;
   } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
     // Explicit extension of mapped data - not allowed.
     MESSAGE("explicit extension not allowed: host address specified is " DPxMOD
             " (%" PRId64 " bytes), but device allocation maps to host at "
             DPxMOD " (%" PRId64 " bytes)",
             DPxPTR(HstPtrBegin), Size, DPxPTR(lr.Entry->HstPtrBegin),
             lr.Entry->HstPtrEnd - lr.Entry->HstPtrBegin);

openmp/libomptarget/src/interface.cpp

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() > 1)
+      if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE)
         for (const auto &Device : PM->Devices)
-          dumpTargetPointerMappings(Device);
+          dumpTargetPointerMappings(loc, Device);
       else
-        FAILURE_MESSAGE("run with env LIBOMPTARGET_INFO>1 to dump host-target "
-                        "pointer maps\n");
+        FAILURE_MESSAGE("Run with LIBOMPTARGET_DEBUG=%d to dump host-target "
+                        "pointer mappings.\n",
+                        OMP_INFOTYPE_DUMP_TABLE);
 
       SourceInfo info(loc);
       if (info.isAvailible())
         fprintf(stderr, "%s:%d:%d: ", info.getFilename(), info.getLine(),
                 info.getColumn());
       else
-        FAILURE_MESSAGE(
-            "Build with debug information to provide more information");
+        FAILURE_MESSAGE("Source location information not present. Compile with "
+                        "-g or -gline-tables-only.\n");
       FATAL_MESSAGE0(
           1, "failure of target construct while offloading is mandatory");
+    } else {
+      if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE)
+        for (const auto &Device : PM->Devices)
+          dumpTargetPointerMappings(loc, Device);
     }
     break;
   }
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 /// adds requires flags
 EXTERN void __tgt_register_requires(int64_t flags) {
EXTERN void __tgt_target_data_begin_mapper(ident_t *loc, int64_t device_id,
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     DP("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false, loc);
     return;
   }
 
   DeviceTy &Device = PM->Devices[device_id];
 
+  if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+    printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+                         arg_names);
 #ifdef OMPTARGET_DEBUG
   for (int i = 0; i < arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
   }
 #endif
EXTERN void __tgt_target_data_end_mapper(ident_t *loc, int64_t device_id,
   }
 
   DeviceTy &Device = PM->Devices[device_id];
   if (!Device.IsInit) {
     DP("Uninit device: ignore");
     HandleTargetOutcome(false, loc);
     return;
   }
 
 #ifdef OMPTARGET_DEBUG
   for (int i=0; i<arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
   }
 #endif
EXTERN void __tgt_target_data_update_mapper(ident_t *loc, int64_t device_id,
   }
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     DP("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false, loc);
     return;
   }
 
+  if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+    printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+                         arg_names);
+
   DeviceTy &Device = PM->Devices[device_id];
   int rc = targetDataUpdate(Device, arg_num, args_base, args, arg_sizes,
                             arg_types, arg_names, arg_mappers);
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
 }
 
 EXTERN void __tgt_target_data_update_nowait_mapper(
     ident_t *loc, int64_t device_id, int32_t arg_num, void **args_base,
EXTERN int __tgt_target_mapper(ident_t *loc, int64_t device_id, void *host_ptr,
   }
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     REPORT("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false, loc);
     return OFFLOAD_FAIL;
   }
 
+  if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+    printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+                         arg_names);
 #ifdef OMPTARGET_DEBUG
   for (int i=0; i<arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
   }
 #endif
EXTERN int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id,
   }
 
   if (CheckDeviceAndCtors(device_id) != OFFLOAD_SUCCESS) {
     REPORT("Failed to get device %" PRId64 " ready\n", device_id);
     HandleTargetOutcome(false, loc);
     return OFFLOAD_FAIL;
   }
 
+  if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
+    printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
+                         arg_names);
 #ifdef OMPTARGET_DEBUG
   for (int i=0; i<arg_num; ++i) {
     DP("Entry %2d: Base=" DPxMOD ", Begin=" DPxMOD ", Size=%" PRId64
        ", Type=0x%" PRIx64 ", Name=%s\n",
        i, DPxPTR(args_base[i]), DPxPTR(args[i]), arg_sizes[i], arg_types[i],
        (arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "(null)");
   }
 #endif
 
   int rc = target(device_id, host_ptr, arg_num, args_base, args, arg_sizes,
                   arg_types, arg_names, arg_mappers, team_num, thread_limit,
                   true /*team*/);
   HandleTargetOutcome(rc == OFFLOAD_SUCCESS, loc);
-
   return rc;
 }
 
 EXTERN int __tgt_target_teams_nowait_mapper(
     ident_t *loc, int64_t device_id, void *host_ptr, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types,
     map_var_info_t *arg_names, void **arg_mappers, int32_t team_num,
     int32_t thread_limit, int32_t depNum, void *depList, int32_t noAliasDepNum,

openmp/libomptarget/src/private.h

 //===---------- private.h - Target independent OpenMP target RTL ----------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Private function declarations and helper macros for debugging output.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_PRIVATE_H
 #define _OMPTARGET_PRIVATE_H
 
-#include <omptarget.h>
 #include <Debug.h>
+#include <SourceInfo.h>
+#include <omptarget.h>
 
 #include <cstdint>
 
 extern int targetDataBegin(DeviceTy &Device, int32_t arg_num, void **args_base,
                            void **args, int64_t *arg_sizes, int64_t *arg_types,
                            map_var_info_t *arg_names, void **arg_mappers,
                            __tgt_async_info *async_info_ptr);
 
 extern int targetDataEnd(DeviceTy &Device, int32_t ArgNum, void **ArgBases,
                          void **Args, int64_t *ArgSizes, int64_t *ArgTypes,
                          map_var_info_t *arg_names, void **ArgMappers,
                          __tgt_async_info *AsyncInfo);
 
 extern int targetDataUpdate(DeviceTy &Device, int32_t arg_num, void **args_base,
                             void **args, int64_t *arg_sizes, int64_t *arg_types,
                             map_var_info_t *arg_names, void **arg_mappers,
                             __tgt_async_info *async_info_ptr = nullptr);
 
 extern int target(int64_t DeviceId, void *HostPtr, int32_t ArgNum,
                   void **ArgBases, void **Args, int64_t *ArgSizes,
                   int64_t *ArgTypes, map_var_info_t *arg_names,
                   void **ArgMappers, int32_t TeamNum, int32_t ThreadLimit,
struct MapComponentInfoTy {
   MapComponentInfoTy(void *Base, void *Begin, int64_t Size, int64_t Type)
       : Base(Base), Begin(Begin), Size(Size), Type(Type) {}
 };
 
 // This structure stores all components of a user-defined mapper. The number of
 // components are dynamically decided, so we utilize C++ STL vector
 // implementation here.
 struct MapperComponentsTy {
   std::vector<MapComponentInfoTy> Components;
   int32_t size() { return Components.size(); }
 };
 
 // The mapper function pointer type. It follows the signature below:
 // void .omp_mapper.<type_name>.<mapper_id>.(void *rt_mapper_handle,
 //                                           void *base, void *begin,
 //                                           size_t size, int64_t type);
 typedef void (*MapperFuncPtrTy)(void *, void *, void *, int64_t, int64_t);
 
 // Function pointer type for target_data_* functions (targetDataBegin,
 // targetDataEnd and targetDataUpdate).
 typedef int (*TargetDataFuncPtrTy)(DeviceTy &, int32_t, void **, void **,
                                    int64_t *, int64_t *, map_var_info_t *,
                                    void **, __tgt_async_info *);
 
 // Implemented in libomp, they are called from within __tgt_* functions.
 #ifdef __cplusplus
 extern "C" {
 #endif
 // functions that extract info from libomp; keep in sync
 int omp_get_default_device(void) __attribute__((weak));
 int32_t __kmpc_omp_taskwait(void *loc_ref, int32_t gtid) __attribute__((weak));
 int32_t __kmpc_global_thread_num(void *) __attribute__((weak));
 int __kmpc_get_target_offload(void) __attribute__((weak));
 #ifdef __cplusplus
 }
 #endif
 
 #define TARGET_NAME Libomptarget
 #define DEBUG_PREFIX GETNAME(TARGET_NAME)
 
 ////////////////////////////////////////////////////////////////////////////////
 /// dump a table of all the host-target pointer pairs on failure
-static inline void dumpTargetPointerMappings(const DeviceTy &Device) {
+static inline void dumpTargetPointerMappings(const ident_t *loc,
+                                             const DeviceTy &Device) {
   if (Device.HostDataToTargetMap.empty())
     return;
 
-  fprintf(stderr, "Device %d Host-Device Pointer Mappings:\n", Device.DeviceID);
-  fprintf(stderr, "%-18s %-18s %s %s\n", "Host Ptr", "Target Ptr", "Size (B)",
-          "Declaration");
+  SourceInfo kernel(loc);
+  INFO((int)Device.DeviceID, "Host-Device Pointer Mappings at %s:%d:%d:\n",
+       kernel.getFilename(), kernel.getLine(), kernel.getColumn());
+  INFO((int)Device.DeviceID, "%-18s %-18s %s %s %s\n", "Host Ptr", "Target Ptr",
+       "Size (B)", "RefCount", "Declaration");
   for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
     SourceInfo info(HostTargetMap.HstPtrName);
-    fprintf(stderr, DPxMOD " " DPxMOD " %-8lu %s at %s:%d:%d\n",
-            DPxPTR(HostTargetMap.HstPtrBegin),
-            DPxPTR(HostTargetMap.TgtPtrBegin),
-            HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin, info.getName(),
-            info.getFilename(), info.getLine(), info.getColumn());
+    INFO((int)Device.DeviceID,
+         DPxMOD " " DPxMOD " %-8lu %-8ld %s at %s:%d:%d\n",
+         DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
+         HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin,
+         HostTargetMap.getRefCount(), info.getName(), info.getFilename(),
+         info.getLine(), info.getColumn());
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// Print out the names and properties of the arguments to each kernel
+static inline void
+printKernelArguments(const ident_t *loc, const int64_t DeviceId,
+                     const int32_t argNum, const int64_t *argSizes,
+                     const int64_t *argTypes, const map_var_info_t *argNames) {
+  SourceInfo info(loc);
+  INFO((int)DeviceId, "Enterting kernel at %s:%d:%d with %d arguments:\n",
+       info.getFilename(), info.getLine(), info.getColumn(), argNum);
+
+  for (int32_t i = 0; i < argNum; ++i) {
+    const map_var_info_t varName = (argNames) ? argNames[i] : nullptr;
+    const char *type = nullptr;
+    const char *implicit =
+        (argTypes[i] & OMP_TGT_MAPTYPE_IMPLICIT) ? "(implicit)" : "";
+    if (argTypes[i] & OMP_TGT_MAPTYPE_TO && argTypes[i] & OMP_TGT_MAPTYPE_FROM)
+      type = "tofrom";
+    else if (argTypes[i] & OMP_TGT_MAPTYPE_TO)
+      type = "to";
+    else if (argTypes[i] & OMP_TGT_MAPTYPE_FROM)
+      type = "from";
+    else if (argTypes[i] & OMP_TGT_MAPTYPE_PRIVATE)
+      type = "private";
+    else if (argTypes[i] & OMP_TGT_MAPTYPE_LITERAL)
+      type = "firstprivate";
+    else if (argTypes[i] & OMP_TGT_MAPTYPE_TARGET_PARAM && argSizes[i] != 0)
+      type = "alloc";
+    else
+      type = "use_address";
+
+    INFO((int)DeviceId, "%s(%s)[%ld] %s\n", type,
+         getNameFromMapping(varName).c_str(), argSizes[i], implicit);
   }
 }
 
 #ifdef OMPTARGET_PROFILE_ENABLED
 #include "llvm/Support/TimeProfiler.h"
 #define TIMESCOPE() llvm::TimeTraceScope TimeScope(__FUNCTION__)
 #else
 #define TIMESCOPE()
 #endif
 
 #endif

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
+// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -gline-tables-only && env LIBOMPTARGET_INFO=23 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
 
 #include <stdio.h>
 #include <omp.h>
 
+#define N 64
+
 int main() {
-    int ptr = 1;
+  int A[N];
+  int B[N];
+  int C[N];
+  int val = 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;}
+// INFO: CUDA device 0 info: Device supports up to {{.*}} CUDA blocks and {{.*}} threads with a warp size of {{.*}}
+// INFO: Libomptarget device 0 info: Enterting kernel at info.c:32:1 with 3 arguments:
+// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
+// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
+// INFO: Libomptarget device 0 info: to(C[0:64])[256]
+// INFO: Libomptarget device 0 info: Host-Device Pointer Mappings at info.c:32:1:
+// INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) RefCount Declaration
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        C[0:64] at info.c:11:7
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        B[0:64] at info.c:10:7
+// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256      1        A[0:64] at info.c:9:7
+// INFO: Libomptarget device 0 info: Enterting kernel at info.c:33:1 with 1 arguments:
+// INFO: Libomptarget device 0 info: firstprivate(val)[4]
+// INFO: CUDA device 0 info: Launching kernel {{.*}} with {{.*}} and {{.*}} threads in {{.*}} mode
+// INFO: Libomptarget device 0 info: Host-Device Pointer Mappings at info.c:33:1:
+// INFO: Libomptarget device 0 info: Host Ptr           Target Ptr         Size (B) RefCount Declaration
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        C[0:64] at info.c:11:7
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        B[0:64] at info.c:10:7
+// INFO: Libomptarget device 0 info: 0x{{.*}} 0x{{.*}} 256      1        A[0:64] at info.c:9:7
+#pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
+#pragma omp target firstprivate(val)
+  { val = 1; }
 
   return 0;
 }