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 | b4efbcf3febda3a201cca9cc8c73f2613c14a828 (patch) | |
| tree | 75ed7423eb92c22d8b8e29f130bbdc9cc05bf0e4 | |
| parent | 1f1c3ec2633a123dc4c88c5b88c7fa60d6867426 (diff) | |
| download | openmp_llvm-b4efbcf3febda3a201cca9cc8c73f2613c14a828.tar.gz | |
[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
git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@350128 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | libomptarget/deviceRTLs/nvptx/src/debug.h | 27 |
1 files changed, 19 insertions, 8 deletions
diff --git a/libomptarget/deviceRTLs/nvptx/src/debug.h b/libomptarget/deviceRTLs/nvptx/src/debug.h index 8577c8f..100713b 100644 --- a/libomptarget/deviceRTLs/nvptx/src/debug.h +++ b/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); \ } \ } |