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 @@ -110,6 +121,7 @@ fprintf(stderr, __VA_ARGS__); \ } +/// Emit a message for debugging #define DP(...) \ do { \ if (getDebugLevel() > 0) { \ @@ -117,6 +129,7 @@ } \ } while (false) +/// Emit a message for debugging or failure if debugging is disabled #define REPORT(...) \ do { \ if (getDebugLevel() > 0) { \ @@ -133,4 +146,14 @@ #define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__); #endif // OMPTARGET_DEBUG +/// Emit a message giving the user extra information about the runtime if +#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); \ @@ -287,6 +287,16 @@ return false; } + const __tgt_offload_entry *getOffloadEntry(const int DeviceId, + const void *Addr) const { + for (const __tgt_offload_entry &Itr : + DeviceData[DeviceId].FuncGblEntries.back().Entries) + if (Itr.addr == Addr) + return &Itr; + + return nullptr; + } + // Return the pointer to the target entries table __tgt_target_table *getOffloadEntriesTable(const int DeviceId) { FuncOrGblEntryTy &E = DeviceData[DeviceId].FuncGblEntries.back(); @@ -492,9 +502,10 @@ 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) { @@ -926,9 +937,14 @@ CudaBlocksPerGrid = TeamNum; } - // Run on the device. - DP("Launch kernel with %d blocks and %d threads\n", CudaBlocksPerGrid, - CudaThreadsPerBlock); + INFO("Device %d launching kernel %s with %d blocks and %d threads in %s " + "mode\n", + DeviceId, + (findOffloadEntry(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, 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 +#include #include #include @@ -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;