Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h
+++ openmp/trunk/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("<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);               \
     }                                                                          \
Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omp_data.cu
+++ openmp/trunk/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
 ////////////////////////////////////////////////////////////////////////////////
 
Index: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
===================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -379,6 +379,19 @@
   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
 ////////////////////////////////////////////////////////////////////////////////
Index: openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp
===================================================================
--- openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp
+++ openmp/trunk/libomptarget/plugins/cuda/src/rtl.cpp
@@ -80,6 +80,12 @@
       : 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 @@
     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);
 }