This is an archive of the discontinued LLVM Phabricator instance.

Differential D55436

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

Authored by ABataev on Dec 7 2018, 7:12 AM.

Details

Reviewers
gtbercea
kkwli0
Commits
rG8acafff404f7: [OPENMP][NVPTX]Save registers for optimized builds with enabled logging.
rOMP348606: [OPENMP][NVPTX]Save registers for optimized builds with enabled logging.
rL348606: [OPENMP][NVPTX]Save registers for optimized builds with enabled logging.
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.

Diff Detail

Repository
rL LLVM

Event Timeline

ABataev created this revision.Dec 7 2018, 7:12 AM
Herald added a subscriber: guansong. · View Herald TranscriptDec 7 2018, 7:12 AM
Harbormaster completed remote builds in B25829: Diff 177211.Dec 7 2018, 7:12 AM
gtbercea accepted this revision.Dec 7 2018, 7:40 AM

LG

This revision is now accepted and ready to land.Dec 7 2018, 7:40 AM
Closed by commit rL348606: [OPENMP][NVPTX]Save registers for optimized builds with enabled logging. (authored by ABataev). · Explain WhyDec 7 2018, 8:11 AM
This revision was automatically updated to reflect the committed changes.
Herald added a subscriber: llvm-commits. · View Herald TranscriptDec 7 2018, 8:11 AM

Revision Contents

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

Diff 177224

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

Show First 20 Lines • Show All 121 Lines • ▼ Show 20 Lines
#endif #endif
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
// implemtation for debug // implemtation for debug
//////////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////
#if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING #if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
#include <stdio.h> #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 #endif
#if OMPTARGET_NVPTX_TEST #if OMPTARGET_NVPTX_TEST
#include <assert.h> #include <assert.h>
#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)
Show All 21 Lines
// implement // implement
#if OMPTARGET_NVPTX_DEBUG #if OMPTARGET_NVPTX_DEBUG
#define DON(_flag) ((unsigned)(OMPTARGET_NVPTX_DEBUG) & (_flag)) #define DON(_flag) ((unsigned)(OMPTARGET_NVPTX_DEBUG) & (_flag))
#define PRINT0(_flag, _str) \ #define PRINT0(_flag, _str) \
{ \ { \
if (omptarget_device_environment.debug_level && DON(_flag)) { \ if (omptarget_device_environment.debug_level && DON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d>: " _str, (int)blockIdx.x, \ log("<b %2d, t %4d, w %2d, l %2d>: " _str); \
(int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
(int)(threadIdx.x & 0x1F)); \
} \ } \
} }
#define PRINT(_flag, _str, _args...) \ #define PRINT(_flag, _str, _args...) \
{ \ { \
if (omptarget_device_environment.debug_level && DON(_flag)) { \ if (omptarget_device_environment.debug_level && DON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d>: " _str, (int)blockIdx.x, \ log("<b %2d, t %4d, w %2d, l %2d>: " _str, _args); \
(int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
(int)(threadIdx.x & 0x1F), _args); \
} \ } \
} }
#else #else
#define DON(_flag) (FALSE) #define DON(_flag) (FALSE)
#define PRINT0(flag, str) #define PRINT0(flag, str)
#define PRINT(flag, str, _args...) #define PRINT(flag, str, _args...)
Show All 27 Lines#define ASSERT(_flag, _cond, _str, _args...) \
} }
#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) && !(_cond)) { \
printf("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n", \ log("<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)); \
assert(_cond); \ assert(_cond); \
} \ } \
} }
#define ASSERT(_flag, _cond, _str, _args...) \ #define ASSERT(_flag, _cond, _str, _args...) \
{ \ { \
if (TON(_flag) && !(_cond)) { \ if (TON(_flag) && !(_cond)) { \
printf("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", \ log("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", _args); \
(int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
(int)(threadIdx.x & 0x1F), _args); \
assert(_cond); \ assert(_cond); \
} \ } \
} }
#else #else
#define TON(_flag) (FALSE) #define TON(_flag) (FALSE)
#define ASSERT0(_flag, _cond, _str) #define ASSERT0(_flag, _cond, _str)
Show All 10 Lines
#define LW_FUSSY (LW_SET_FUSSY) #define LW_FUSSY (LW_SET_FUSSY)
#if OMPTARGET_NVPTX_WARNING #if OMPTARGET_NVPTX_WARNING
#define WON(_flag) ((OMPTARGET_NVPTX_WARNING) & (_flag)) #define WON(_flag) ((OMPTARGET_NVPTX_WARNING) & (_flag))
#define WARNING0(_flag, _str) \ #define WARNING0(_flag, _str) \
{ \ { \
if (WON(_flag)) { \ if (WON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, (int)blockIdx.x, \ log("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str); \
(int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
(int)(threadIdx.x & 0x1F)); \
} \ } \
} }
#define WARNING(_flag, _str, _args...) \ #define WARNING(_flag, _str, _args...) \
{ \ { \
if (WON(_flag)) { \ if (WON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, (int)blockIdx.x, \ log("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, _args); \
(int)threadIdx.x, (int)(threadIdx.x / WARPSIZE), \
(int)(threadIdx.x & 0x1F), _args); \
} \ } \
} }
#else #else
#define WON(_flag) (FALSE) #define WON(_flag) (FALSE)
#define WARNING0(_flag, _str) #define WARNING0(_flag, _str)
#define WARNING(_flag, _str, _args...) #define WARNING(_flag, _str, _args...)
#endif #endif
#endif #endif