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 @@ -100,6 +100,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 @@ -133,4 +139,13 @@ #define REPORT(...) FAILURE_MESSAGE(__VA_ARGS__); #endif // OMPTARGET_DEBUG +#define INFO(...) \ + do { \ + if (getInfoLevel() > 0) { \ + DP(__VA_ARGS__); \ + } else { \ + 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 @@ -17,12 +17,27 @@ #include #include +#include #include // Store target policy (disabled, mandatory, default) kmp_target_offload_kind_t TargetOffloadPolicy = tgt_default; std::mutex TargetOffloadMtx; +//////////////////////////////////////////////////////////////////////////////// +/// manage the success or failure of a target construct +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;