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 | 956e323841a346d80d6b9bbacebb0f3cecf92cd9 (patch) | |
| tree | 744b5c593cd9d17cd66ef5c6f4b5abe60be04f58 /openmp/libomptarget | |
| parent | ef46b214425df92aaa7dfee2b1627cb4970221f7 (diff) | |
[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
Diffstat (limited to 'openmp/libomptarget')
| -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 b556670f182..8577c8f9599 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); \ } \ } |