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
@@ -127,6 +127,14 @@
 
 #if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
 #include <stdio.h>
+#include "option.h"
+
+template <typename... Arguments>
+static NOINLINE void log(const char *fmt, Arguments... parameters) {
+  printf(fmt, (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),
+         (int)(threadIdx.x & 0x1F), parameters...);
+}
+
 #endif
 #if OMPTARGET_NVPTX_TEST
 #include <assert.h>
@@ -164,18 +172,14 @@
 #define PRINT0(_flag, _str)                                                    \
   {                                                                            \
     if (omptarget_device_environment.debug_level && DON(_flag)) {              \
-      printf("<b %2d, t %4d, w %2d, l %2d>: " _str, (int)blockIdx.x,           \
-             (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),                  \
-             (int)(threadIdx.x & 0x1F));                                       \
+      log("<b %2d, t %4d, w %2d, l %2d>: " _str);                              \
     }                                                                          \
   }
 
 #define PRINT(_flag, _str, _args...)                                           \
   {                                                                            \
     if (omptarget_device_environment.debug_level && DON(_flag)) {              \
-      printf("<b %2d, t %4d, w %2d, l %2d>: " _str, (int)blockIdx.x,           \
-             (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),                  \
-             (int)(threadIdx.x & 0x1F), _args);                                \
+      log("<b %2d, t %4d, w %2d, l %2d>: " _str, _args);                       \
     }                                                                          \
   }
 #else
@@ -219,18 +223,14 @@
 #define ASSERT0(_flag, _cond, _str)                                            \
   {                                                                            \
     if (TON(_flag) && !(_cond)) {                                              \
-      printf("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n",                \
-             (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
-             (int)(threadIdx.x & 0x1F));                                       \
+      log("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n");                  \
       assert(_cond);                                                           \
     }                                                                          \
   }
 #define ASSERT(_flag, _cond, _str, _args...)                                   \
   {                                                                            \
     if (TON(_flag) && !(_cond)) {                                              \
-      printf("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n",                \
-             (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
-             (int)(threadIdx.x & 0x1F), _args);                                \
+      log("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", _args);           \
       assert(_cond);                                                           \
     }                                                                          \
   }
@@ -257,17 +257,13 @@
 #define WARNING0(_flag, _str)                                                  \
   {                                                                            \
     if (WON(_flag)) {                                                          \
-      printf("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, (int)blockIdx.x,   \
-             (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),                  \
-             (int)(threadIdx.x & 0x1F));                                       \
+      log("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str);                      \
     }                                                                          \
   }
 #define WARNING(_flag, _str, _args...)                                         \
   {                                                                            \
     if (WON(_flag)) {                                                          \
-      printf("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, (int)blockIdx.x,   \
-             (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),                  \
-             (int)(threadIdx.x & 0x1F), _args);                                \
+      log("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, _args);               \
     }                                                                          \
   }