Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2024-10-07 19:03:57 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[OpenMP] Use LIBOMPTARGET_DEVICE_RTL_DEBUG env var to control debug messages on the device side

Browse Source 
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
This commit is contained in:
Guansong Zhang 2018-05-04 19:29:28 +00:00
parent 8e4958e760
commit e1c7a46d5b
4 changed files with 69 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
View File

@ -163,7 +163,7 @@
#define PRINT0(_flag, _str) \
{ \
if (DON(_flag)) { \
if (omptarget_device_environment.debug_level && DON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d>: " _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("<b %2d, t %4d, w %2d, l %2d>: " _str, blockIdx.x, threadIdx.x, \
threadIdx.x / WARPSIZE, threadIdx.x & 0x1F, _args); \
} \

6
openmp/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
View File

@ -13,6 +13,12 @@
#include "omptarget-nvptx.h"
////////////////////////////////////////////////////////////////////////////////
// global device envrionment
////////////////////////////////////////////////////////////////////////////////
__device__ omptarget_device_environmentTy omptarget_device_environment;
////////////////////////////////////////////////////////////////////////////////
// global data holding OpenMP state information
////////////////////////////////////////////////////////////////////////////////

13
openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
View File

@ -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
////////////////////////////////////////////////////////////////////////////////

48
openmp/libomptarget/plugins/cuda/src/rtl.cpp
View File

@ -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<KernelTy> 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);
}