diff options
| author | Alexey Bataev <a.bataev@hotmail.com> | 2018-12-07 16:08:29 +0000 |
|---|---|---|
| committer | Alexey Bataev <a.bataev@hotmail.com> | 2018-12-07 16:08:29 +0000 |
| commit | 8acafff404f7b8f18d278ad4caadf86b0b37995d (patch) | |
| tree | 8e6eb03a78f5b261a41ea3ce87f3cb995c577f2c /openmp | |
| parent | b87251d0bb9309113f234e020720d3505938e1b1 (diff) | |
| download | llvm-8acafff404f7b8f18d278ad4caadf86b0b37995d.zip llvm-8acafff404f7b8f18d278ad4caadf86b0b37995d.tar.gz llvm-8acafff404f7b8f18d278ad4caadf86b0b37995d.tar.bz2 | |
[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
llvm-svn: 348606
Diffstat (limited to 'openmp')
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/debug.h | 32 |
1 files changed, 14 insertions, 18 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h index b556670..8577c8f 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h +++ b/openmp/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); \ } \ } |