[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
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@348606 91177308-0d34-0410-b5e6-96231b3b80d8
diff --git a/libomptarget/deviceRTLs/nvptx/src/debug.h b/libomptarget/deviceRTLs/nvptx/src/debug.h
index b556670..8577c8f 100644
--- a/libomptarget/deviceRTLs/nvptx/src/debug.h
+++ b/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); \
} \
}
