diff options
| author | Alexey Bataev <a.bataev@hotmail.com> | 2018-12-28 17:29:47 +0000 |
|---|---|---|
| committer | Alexey Bataev <a.bataev@hotmail.com> | 2018-12-28 17:29:47 +0000 |
| commit | a6eb216e054b3d6a008534d263d4bf2bbad9bb52 (patch) | |
| tree | df802a131c88c8ec90f51d7b6d170774c0eeb53c /openmp | |
| parent | 1fc1c50ce76b7d20cec55fde3a4476d626344b50 (diff) | |
[OPENMP][NVPTX]Outline assert into noinline function, NFC.
Summary:
At high optimization level asserts lead to some unexpected results
because of auto-inserted unreachable instructions. This outlining
prevents some of such dangerous optimizations and leads to better
stability.
Reviewers: gtbercea, kkwli0, grokos
Subscribers: guansong, caomhin, openmp-commits
Differential Revision: https://reviews.llvm.org/D56101
Diffstat (limited to 'openmp')
| -rw-r--r-- | openmp/libomptarget/deviceRTLs/nvptx/src/debug.h | 27 |
1 files changed, 19 insertions, 8 deletions
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h index 8577c8f9599..100713bbca0 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/debug.h @@ -138,6 +138,18 @@ static NOINLINE void log(const char *fmt, Arguments... parameters) { #endif #if OMPTARGET_NVPTX_TEST #include <assert.h> + +template <typename... Arguments> +NOINLINE static void check(bool cond, const char *fmt, + Arguments... parameters) { + if (!cond) + printf(fmt, (int)blockIdx.x, (int)threadIdx.x, + (int)(threadIdx.x / WARPSIZE), (int)(threadIdx.x & 0x1F), + parameters...); + assert(cond); +} + +NOINLINE static void check(bool cond) { assert(cond); } #endif // set flags that are tested (inclusion properties) @@ -207,13 +219,13 @@ static NOINLINE void log(const char *fmt, Arguments... parameters) { #define ASSERT0(_flag, _cond, _str) \ { \ if (TON(_flag)) { \ - assert(_cond); \ + check(_cond); \ } \ } #define ASSERT(_flag, _cond, _str, _args...) \ { \ if (TON(_flag)) { \ - assert(_cond); \ + check(_cond); \ } \ } @@ -222,16 +234,15 @@ static NOINLINE void log(const char *fmt, Arguments... parameters) { #define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag)) #define ASSERT0(_flag, _cond, _str) \ { \ - if (TON(_flag) && !(_cond)) { \ - log("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n"); \ - assert(_cond); \ + if (TON(_flag)) { \ + check((_cond), "<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n"); \ } \ } #define ASSERT(_flag, _cond, _str, _args...) \ { \ - if (TON(_flag) && !(_cond)) { \ - log("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", _args); \ - assert(_cond); \ + if (TON(_flag)) { \ + check((_cond), "<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", \ + _args); \ } \ } |