[PATCH] D55436: [OPENMP][NVPTX]Save registers for optimized builds with enabled logging.

[Alexey Bataev via Phabricator via llvm-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rL348606: [OPENMP][NVPTX]Save registers for optimized builds with enabled logging. (authored by ABataev, committed by ).
Herald added a subscriber: llvm-commits.

Repository:
  rL LLVM

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D55436/new/

https://reviews.llvm.org/D55436

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


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);               \
     }                                                                          \
   }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D55436.177224.patch
Type: text/x-patch
Size: 4897 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20181207/8102b038/attachment.bin>