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

[Alexey Bataev via Phabricator via Openmp-commits]

ABataev created this revision.
ABataev added reviewers: gtbercea, kkwli0.
Herald added a subscriber: guansong.

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.


Repository:
  rOMP OpenMP

https://reviews.llvm.org/D55436

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


Index: libomptarget/deviceRTLs/nvptx/src/debug.h
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/debug.h
+++ 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.177211.patch
Type: text/x-patch
Size: 4858 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20181207/d0ca5fb0/attachment.bin>