intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-02-08 17:28:30 +08:00
Code Issues Packages Projects Releases Wiki Activity

[OPENMP][NVPTX]Outline assert into noinline function, NFC.

Browse Source 
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

llvm-svn: 350128
This commit is contained in:
Alexey Bataev
2018-12-28 17:29:47 +00:00
parent db43f0696e
commit 1708858dbd
1 changed files with 19 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

27
openmp/libomptarget/deviceRTLs/nvptx/src/debug.h
View File

@@ -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); \
} \
}