This is an archive of the discontinued LLVM Phabricator instance.

Differential D56101

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

Authored by ABataev on Dec 27 2018, 7:01 AM.

Details

Reviewers
gtbercea
kkwli0
grokos
Commits
rG1708858dbd89: [OPENMP][NVPTX]Outline assert into noinline function, NFC.
rOMP350128: [OPENMP][NVPTX]Outline assert into noinline function, NFC.
rL350128: [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.

Diff Detail

Repository
rL LLVM

Event Timeline

ABataev created this revision.Dec 27 2018, 7:01 AM
Herald added a subscriber: guansong. · View Herald TranscriptDec 27 2018, 7:01 AM
Harbormaster completed remote builds in B26273: Diff 179550.Dec 27 2018, 7:03 AM
grokos accepted this revision.Dec 27 2018, 8:15 AM

Looks good.

This revision is now accepted and ready to land.Dec 27 2018, 8:15 AM
Closed by commit rL350128: [OPENMP][NVPTX]Outline assert into noinline function, NFC. (authored by ABataev). · Explain WhyDec 28 2018, 9:33 AM
This revision was automatically updated to reflect the committed changes.
Herald added a subscriber: llvm-commits. · View Herald TranscriptDec 28 2018, 9:33 AM

Revision Contents

PathSize
openmp/
trunk/
libomptarget/
deviceRTLs/
nvptx/
src/
27 lines

Diff 179645

openmp/trunk/libomptarget/deviceRTLs/nvptx/src/debug.h

Show First 20 Lines • Show All 132 Lines • ▼ Show 20 Lines
static NOINLINE void log(const char *fmt, Arguments... parameters) { static NOINLINE void log(const char *fmt, Arguments... parameters) {
printf(fmt, (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), printf(fmt, (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),
(int)(threadIdx.x & 0x1F), parameters...); (int)(threadIdx.x & 0x1F), parameters...);
} }
#endif #endif
#if OMPTARGET_NVPTX_TEST #if OMPTARGET_NVPTX_TEST
#include <assert.h> #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 #endif
// set flags that are tested (inclusion properties) // set flags that are tested (inclusion properties)
#define LD_ALL (LD_SET_ALL) #define LD_ALL (LD_SET_ALL)
#define LD_LOOP (LD_SET_LOOP | LD_SET_LOOPD) #define LD_LOOP (LD_SET_LOOP | LD_SET_LOOPD)
#define LD_LOOPD (LD_SET_LOOPD) #define LD_LOOPD (LD_SET_LOOPD)
▲ Show 20 Lines • Show All 53 Lines • ▼ Show 20 Lines
#define LT_FUSSY (LT_SET_FUSSY) #define LT_FUSSY (LT_SET_FUSSY)
#if OMPTARGET_NVPTX_TEST == LT_SET_SAFETY #if OMPTARGET_NVPTX_TEST == LT_SET_SAFETY
#define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag)) #define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag))
#define ASSERT0(_flag, _cond, _str) \ #define ASSERT0(_flag, _cond, _str) \
{ \ { \
if (TON(_flag)) { \ if (TON(_flag)) { \
assert(_cond); \ check(_cond); \
} \ } \
} }
#define ASSERT(_flag, _cond, _str, _args...) \ #define ASSERT(_flag, _cond, _str, _args...) \
{ \ { \
if (TON(_flag)) { \ if (TON(_flag)) { \
assert(_cond); \ check(_cond); \
} \ } \
} }
#elif OMPTARGET_NVPTX_TEST >= LT_SET_INPUT #elif OMPTARGET_NVPTX_TEST >= LT_SET_INPUT
#define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag)) #define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag))
#define ASSERT0(_flag, _cond, _str) \ #define ASSERT0(_flag, _cond, _str) \
{ \ { \
if (TON(_flag) && !(_cond)) { \ if (TON(_flag)) { \
log("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n"); \ check((_cond), "<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n"); \
assert(_cond); \
} \ } \
} }
#define ASSERT(_flag, _cond, _str, _args...) \ #define ASSERT(_flag, _cond, _str, _args...) \
{ \ { \
if (TON(_flag) && !(_cond)) { \ if (TON(_flag)) { \
log("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", _args); \ check((_cond), "<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", \
assert(_cond); \ _args); \
} \ } \
} }
#else #else
#define TON(_flag) (FALSE) #define TON(_flag) (FALSE)
#define ASSERT0(_flag, _cond, _str) #define ASSERT0(_flag, _cond, _str)
#define ASSERT(_flag, _cond, _str, _args...) #define ASSERT(_flag, _cond, _str, _args...)
Show All 36 Lines