Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2024-11-28 16:11:29 +00:00
Code Issues Actions 2 Packages Projects Releases Wiki Activity

[OPENMP][NVPTX]Save registers for optimized builds with enabled logging.

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

llvm-svn: 348606
This commit is contained in:
Alexey Bataev 2018-12-07 16:08:29 +00:00
parent b87251d0bb
commit 8acafff404
1 changed files with 14 additions and 18 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

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