From e1c7a46d5b55bbe870c0b9aa3a4f3d3ea396550e Mon Sep 17 00:00:00 2001 From: Guansong Zhang Date: Fri, 4 May 2018 19:29:28 +0000 Subject: [PATCH] [OpenMP] Use LIBOMPTARGET_DEVICE_RTL_DEBUG env var to control debug messages on the device side Summary: Enable the device side debug messages at compile time, use env var to control at runtime. To achieve this, an environment data block is passed to the device lib when it is loaded. By default, the message is off, to enable it, a user need to set LIBOMPDEVICE_DEBUG=1. Reviewers: grokos Reviewed By: grokos Subscribers: openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D46210 llvm-svn: 331550 --- openmp/libomptarget/deviceRTLs/nvptx/src/debug.h | 4 +- .../libomptarget/deviceRTLs/nvptx/src/omp_data.cu | 6 +++ .../deviceRTLs/nvptx/src/omptarget-nvptx.h | 13 ++++++ openmp/libomptarget/plugins/cuda/src/rtl.cpp | 48 ++++++++++++++++++++++ 4 files changed, 69 insertions(+), 2 deletions(-) diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h index 1468757..9f59d66 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h @@ -163,7 +163,7 @@ #define PRINT0(_flag, _str) \ { \ - if (DON(_flag)) { \ + if (omptarget_device_environment.debug_level && DON(_flag)) { \ printf(": " _str, blockIdx.x, threadIdx.x, \ threadIdx.x / WARPSIZE, threadIdx.x & 0x1F); \ } \ @@ -171,7 +171,7 @@ #define PRINT(_flag, _str, _args...) \ { \ - if (DON(_flag)) { \ + if (omptarget_device_environment.debug_level && DON(_flag)) { \ printf(": " _str, blockIdx.x, threadIdx.x, \ threadIdx.x / WARPSIZE, threadIdx.x & 0x1F, _args); \ } \ diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu index 33303e7..149af8d 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu @@ -14,6 +14,12 @@ #include "omptarget-nvptx.h" //////////////////////////////////////////////////////////////////////////////// +// global device envrionment +//////////////////////////////////////////////////////////////////////////////// + +__device__ omptarget_device_environmentTy omptarget_device_environment; + +//////////////////////////////////////////////////////////////////////////////// // global data holding OpenMP state information //////////////////////////////////////////////////////////////////////////////// diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h index 0d635a3..2bc5819 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h @@ -379,6 +379,19 @@ private: uint64_t SourceQueue; }; +/// Device envrionment data +struct omptarget_device_environmentTy { + int32_t debug_level; +}; + +//////////////////////////////////////////////////////////////////////////////// +// global device envrionment +//////////////////////////////////////////////////////////////////////////////// + +extern __device__ omptarget_device_environmentTy omptarget_device_environment; + +//////////////////////////////////////////////////////////////////////////////// + //////////////////////////////////////////////////////////////////////////////// // global data tables //////////////////////////////////////////////////////////////////////////////// diff --git a/openmp/libomptarget/plugins/cuda/src/rtl.cpp b/openmp/libomptarget/plugins/cuda/src/rtl.cpp index 920bc4e..fe2f9f6 100644 --- a/openmp/libomptarget/plugins/cuda/src/rtl.cpp +++ b/openmp/libomptarget/plugins/cuda/src/rtl.cpp @@ -80,6 +80,12 @@ struct KernelTy { : Func(_Func), ExecutionMode(_ExecutionMode) {} }; +/// Device envrionment data +/// Manually sync with the deviceRTL side for now, move to a dedicated header file later. +struct omptarget_device_environmentTy { + int32_t debug_level; +}; + /// List that contains all the kernels. /// FIXME: we may need this to be per device and per library. std::list KernelsList; @@ -486,6 +492,48 @@ __tgt_target_table *__tgt_rtl_load_binary(int32_t device_id, DeviceInfo.addOffloadEntry(device_id, entry); } + // send device environment data to the device + { + omptarget_device_environmentTy device_env; + + device_env.debug_level = 0; + +#ifdef OMPTARGET_DEBUG + if (char *envStr = getenv("LIBOMPTARGET_DEVICE_RTL_DEBUG")) { + device_env.debug_level = std::stoi(envStr); + } +#endif + + const char * device_env_Name="omptarget_device_environment"; + CUdeviceptr device_env_Ptr; + size_t cusize; + + err = cuModuleGetGlobal(&device_env_Ptr, &cusize, cumod, device_env_Name); + + if (err == CUDA_SUCCESS) { + if ((size_t)cusize != sizeof(device_env)) { + DP("Global device_environment '%s' - size mismatch (%zu != %zu)\n", + device_env_Name, cusize, sizeof(int32_t)); + CUDA_ERR_STRING(err); + return NULL; + } + + err = cuMemcpyHtoD(device_env_Ptr, &device_env, cusize); + if (err != CUDA_SUCCESS) { + DP("Error when copying data from host to device. Pointers: " + "host = " DPxMOD ", device = " DPxMOD ", size = %zu\n", + DPxPTR(&device_env), DPxPTR(device_env_Ptr), cusize); + CUDA_ERR_STRING(err); + return NULL; + } + + DP("Sending global device environment data %zu bytes\n", (size_t)cusize); + } else { + DP("Finding global device environment '%s' - symbol missing.\n", device_env_Name); + DP("Continue, considering this is a device RTL which does not accept envrionment setting.\n"); + } + } + return DeviceInfo.getOffloadEntriesTable(device_id); } -- 2.7.4