[Openmp-commits] [openmp] r348606 - [OPENMP][NVPTX]Save registers for optimized builds with enabled logging.

[Alexey Bataev via Openmp-commits]

Author: abataev
Date: Fri Dec  7 08:08:29 2018
New Revision: 348606

URL: http://llvm.org/viewvc/llvm-project?rev=348606&view=rev
Log:
[OPENMP][NVPTX]Save registers for optimized builds with enabled logging.

Summary:
Introduced special noinline function log that allows to save some
registers for optimized builds but with enabled logging. Also, it
increases the stability of the optimized builds with inlined runtime.

Reviewers: gtbercea, kkwli0

Reviewed By: gtbercea

Subscribers: caomhin, guansong, openmp-commits

Differential Revision: https://reviews.llvm.org/D55436

Modified:
    openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h

Modified: openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h?rev=348606&r1=348605&r2=348606&view=diff
==============================================================================
--- openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h (original)
+++ openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h Fri Dec  7 08:08:29 2018
@@ -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);               \
     }                                                                          \
   }