Skip to content

Commit

Permalink
[OpenMP] Add using bit flags to select Libomptarget Information
Browse files Browse the repository at this point in the history 
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.

Reviewers: jdoerfert JonChesterfield grokos

Differential Revision: https://reviews.llvm.org/D93727
  • Loading branch information
@jhuber6 @jhuber-ornl
jhuber6 authored and jhuber-ornl committed Jan 4, 2021
1 parent f7463ca commit fe5d51a
Show file tree
Hide file tree
Showing 7 changed files with 177 additions and 67 deletions.
44 changes: 29 additions & 15 deletions openmp/libomptarget/include/Debug.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,24 +37,38 @@
#ifndef _OMPTARGET_DEBUG_H
#define _OMPTARGET_DEBUG_H

static inline int getInfoLevel() {
static int InfoLevel = -1;
if (InfoLevel >= 0)
return InfoLevel;

if (char *EnvStr = getenv("LIBOMPTARGET_INFO"))
InfoLevel = std::stoi(EnvStr);
#include <mutex>

/// 32-Bit field data attributes controlling information presented to the user.
enum OpenMPInfoType : uint32_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 uint32_t getInfoLevel() {
static uint32_t InfoLevel = 0;
static std::once_flag Flag{};
std::call_once(Flag, []() {
if (char *EnvStr = getenv("LIBOMPTARGET_INFO"))
InfoLevel = std::stoi(EnvStr);
});

return InfoLevel;
}

static inline int getDebugLevel() {
static int DebugLevel = -1;
if (DebugLevel >= 0)
return DebugLevel;

if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))
DebugLevel = std::stoi(EnvStr);
static inline uint32_t getDebugLevel() {
static uint32_t DebugLevel = 0;
static std::once_flag Flag{};
std::call_once(Flag, []() {
if (char *EnvStr = getenv("LIBOMPTARGET_DEBUG"))
DebugLevel = std::stoi(EnvStr);
});

return DebugLevel;
}
Expand Down Expand Up @@ -107,7 +121,7 @@ static inline int getDebugLevel() {
/// 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)

Expand Down
9 changes: 8 additions & 1 deletion openmp/libomptarget/include/SourceInfo.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -54,6 +54,13 @@ class SourceInfo {
return std::string(reinterpret_cast<const char *>(name));
}

std::string initStr(const ident_t *loc) {
if (!loc)
return ";unknown;unknown;0;0;;";
else
return std::string(reinterpret_cast<const char *>(loc->psource));
}

/// Get n-th substring in an expression separated by ;.
std::string getSubstring(const int n) const {
std::size_t begin = sourceStr.find(';');
Expand All @@ -73,7 +80,7 @@ class SourceInfo {

public:
SourceInfo(const ident_t *loc)
: sourceStr(initStr(loc->psource)), name(getSubstring(1)),
: sourceStr(initStr(loc)), name(getSubstring(1)),
filename(removePath(getSubstring(0))), line(std::stoi(getSubstring(2))),
column(std::stoi(getSubstring(3))) {}

Expand Down
28 changes: 15 additions & 13 deletions openmp/libomptarget/plugins/cuda/src/rtl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -501,11 +501,12 @@ class DeviceRTLTy {
DeviceData[DeviceId].BlocksPerGrid = EnvTeamLimit;
}

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);
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) {
Expand Down Expand Up @@ -937,14 +938,15 @@ class DeviceRTLTy {
CudaBlocksPerGrid = TeamNum;
}

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");
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,
Expand Down
23 changes: 13 additions & 10 deletions openmp/libomptarget/src/device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -49,10 +49,11 @@ DeviceTy::DeviceTy(RTLInfoTy *RTL)
MemoryManager(nullptr) {}

DeviceTy::~DeviceTy() {
if (DeviceID == -1 || getInfoLevel() < 1)
if (DeviceID == -1 || !(getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE))
return;

dumpTargetPointerMappings(*this);
ident_t loc = {0, 0, 0, 0, ";libomptarget;libomptarget;0;0;;"};
dumpTargetPointerMappings(&loc, *this);
}

int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
Expand Down Expand Up @@ -217,14 +218,16 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
HT.incRefCount();

uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
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(),
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "(null)");
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(),
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
rc = (void *)tp;
} else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
// Explicit extension of mapped data - not allowed.
Expand Down
42 changes: 31 additions & 11 deletions openmp/libomptarget/src/interface.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -57,22 +57,27 @@ static void HandleTargetOutcome(bool success, ident_t *loc = nullptr) {
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;
}
Expand Down Expand Up @@ -147,12 +152,15 @@ EXTERN void __tgt_target_data_begin_mapper(ident_t *loc, int64_t device_id,

DeviceTy &Device = PM->Devices[device_id];

if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
arg_names, "Entering OpenMP data region");
#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)");
(arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif

Expand Down Expand Up @@ -227,12 +235,15 @@ EXTERN void __tgt_target_data_end_mapper(ident_t *loc, int64_t device_id,
return;
}

if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
arg_names, "Exiting OpenMP data region");
#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)");
(arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif

Expand Down Expand Up @@ -294,6 +305,10 @@ EXTERN void __tgt_target_data_update_mapper(ident_t *loc, int64_t device_id,
return;
}

if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
arg_names, "Updating OpenMP data");

DeviceTy &Device = PM->Devices[device_id];
int rc = targetDataUpdate(Device, arg_num, args_base, args, arg_sizes,
arg_types, arg_names, arg_mappers);
Expand Down Expand Up @@ -351,12 +366,15 @@ EXTERN int __tgt_target_mapper(ident_t *loc, int64_t device_id, void *host_ptr,
return OFFLOAD_FAIL;
}

if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
arg_names, "Entering OpenMP kernel");
#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)");
(arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#endif

Expand Down Expand Up @@ -423,20 +441,22 @@ EXTERN int __tgt_target_teams_mapper(ident_t *loc, int64_t device_id,
return OFFLOAD_FAIL;
}

if (getInfoLevel() & OMP_INFOTYPE_KERNEL_ARGS)
printKernelArguments(loc, device_id, arg_num, arg_sizes, arg_types,
arg_names, "Entering OpenMP kernel");
#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)");
(arg_names) ? getNameFromMapping(arg_names[i]).c_str() : "unknown");
}
#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;
}

Expand Down
63 changes: 52 additions & 11 deletions openmp/libomptarget/src/private.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,8 +13,9 @@
#ifndef _OMPTARGET_PRIVATE_H
#define _OMPTARGET_PRIVATE_H

#include <omptarget.h>
#include <Debug.h>
#include <SourceInfo.h>
#include <omptarget.h>

#include <cstdint>

Expand Down Expand Up @@ -90,20 +91,60 @@ int __kmpc_get_target_offload(void) __attribute__((weak));

////////////////////////////////////////////////////////////////////////////////
/// 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(Device.DeviceID,
"OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
INFO(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());
SourceInfo Info(HostTargetMap.HstPtrName);
INFO(Device.DeviceID, DPxMOD " " DPxMOD " %-8lu %-8ld %s at %s:%d:%d\n",
DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
(long unsigned)(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,
const char *RegionType) {
SourceInfo info(Loc);
INFO(DeviceId, "%s at %s:%d:%d with %d arguments:\n", RegionType,
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(DeviceId, "%s(%s)[%ld] %s\n", type,
getNameFromMapping(varName).c_str(), ArgSizes[i], implicit);
}
}

Expand Down
Loading

0 comments on commit fe5d51a

Please sign in to comment.