1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
|
//===------------- debug.h - NVPTX OpenMP debug macros ----------- CUDA -*-===//
//
// The LLVM Compiler Infrastructure
//
// This file is dual licensed under the MIT and the University of Illinois Open
// Source Licenses. See LICENSE.txt for details.
//
//===----------------------------------------------------------------------===//
//
// This file contains debug macros to be used in the application.
//
// Usage guide
//
// PRINT0(flag, str) : if debug flag is on, print (no arguments)
// PRINT(flag, str, args) : if debug flag is on, print (arguments)
// DON(flag) : return true if debug flag is on
//
// ASSERT(flag, cond, str, args): if test flag is on, test the condition
// if the condition is false, print str+args
// and assert.
// CAUTION: cond may be evaluate twice
// AON(flag) : return true if test flag is on
//
// WARNING(flag, str, args) : if warning flag is on, print the warning
// WON(flag) : return true if warning flag is on
//
//===----------------------------------------------------------------------===//
#ifndef _OMPTARGET_NVPTX_DEBUG_H_
#define _OMPTARGET_NVPTX_DEBUG_H_
////////////////////////////////////////////////////////////////////////////////
// set desired level of debugging
////////////////////////////////////////////////////////////////////////////////
#define LD_SET_NONE 0ULL /* none */
#define LD_SET_ALL -1ULL /* all */
// pos 1
#define LD_SET_LOOP 0x1ULL /* basic loop */
#define LD_SET_LOOPD 0x2ULL /* basic loop */
#define LD_SET_PAR 0x4ULL /* basic parallel */
#define LD_SET_PARD 0x8ULL /* basic parallel */
// pos 2
#define LD_SET_SYNC 0x10ULL /* sync info */
#define LD_SET_SYNCD 0x20ULL /* sync info */
#define LD_SET_WAIT 0x40ULL /* state when waiting */
#define LD_SET_TASK 0x80ULL /* print task info (high level) */
// pos 3
#define LD_SET_IO 0x100ULL /* big region io (excl atomic) */
#define LD_SET_IOD 0x200ULL /* big region io (excl atomic) */
#define LD_SET_ENV 0x400ULL /* env info */
#define LD_SET_CANCEL 0x800ULL /* print cancel info */
// pos 4
#define LD_SET_MEM 0x1000ULL /* malloc / free */
////////////////////////////////////////////////////////////////////////////////
// set the desired flags to print selected output.
// these are some examples of possible definitions that can be used for
// debugging.
//#define OMPTARGET_NVPTX_DEBUG (LD_SET_ALL)
//#define OMPTARGET_NVPTX_DEBUG (LD_SET_LOOP) // limit to loop printfs to save
// on cuda buffer
//#define OMPTARGET_NVPTX_DEBUG (LD_SET_IO)
//#define OMPTARGET_NVPTX_DEBUG (LD_SET_IO | LD_SET_ENV)
//#define OMPTARGET_NVPTX_DEBUG (LD_SET_PAR)
#ifndef OMPTARGET_NVPTX_DEBUG
#define OMPTARGET_NVPTX_DEBUG LD_SET_NONE
#elif OMPTARGET_NVPTX_DEBUG
#warning debug is used, not good for measurements
#endif
////////////////////////////////////////////////////////////////////////////////
// set desired level of asserts
////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////////////////////////
// available flags
#define LT_SET_NONE 0x0 /* unsafe */
#define LT_SET_SAFETY \
0x1 /* check malloc type of stuff, input at creation, cheap */
#define LT_SET_INPUT 0x2 /* check also all runtime inputs */
#define LT_SET_FUSSY 0x4 /* fussy checks, expensive */
////////////////////////////////////////////////////////////////////////////////
// set the desired flags
#ifndef OMPTARGET_NVPTX_TEST
#if OMPTARGET_NVPTX_DEBUG
#define OMPTARGET_NVPTX_TEST (LT_SET_FUSSY)
#else
#define OMPTARGET_NVPTX_TEST (LT_SET_SAFETY)
#endif
#endif
////////////////////////////////////////////////////////////////////////////////
// set desired level of warnings
////////////////////////////////////////////////////////////////////////////////
////////////////////////////////////////////////////////////////////////////////
// available flags
#define LW_SET_ALL -1
#define LW_SET_NONE 0x0
#define LW_SET_ENV 0x1
#define LW_SET_INPUT 0x2
#define LW_SET_FUSSY 0x4
////////////////////////////////////////////////////////////////////////////////
// set the desired flags
#if OMPTARGET_NVPTX_DEBUG
#define OMPTARGET_NVPTX_WARNING (LW_SET_NONE)
#else
#define OMPTARGET_NVPTX_WARNING (LW_SET_FUSSY)
#endif
////////////////////////////////////////////////////////////////////////////////
// implemtation for debug
////////////////////////////////////////////////////////////////////////////////
#if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
#include <stdio.h>
#endif
#if OMPTARGET_NVPTX_TEST
#include <assert.h>
#endif
// set flags that are tested (inclusion properties)
#define LD_ALL (LD_SET_ALL)
#define LD_LOOP (LD_SET_LOOP | LD_SET_LOOPD)
#define LD_LOOPD (LD_SET_LOOPD)
#define LD_PAR (LD_SET_PAR | LD_SET_PARD)
#define LD_PARD (LD_SET_PARD)
// pos 2
#define LD_SYNC (LD_SET_SYNC | LD_SET_SYNCD)
#define LD_SYNCD (LD_SET_SYNCD)
#define LD_WAIT (LD_SET_WAIT)
#define LD_TASK (LD_SET_TASK)
// pos 3
#define LD_IO (LD_SET_IO | LD_SET_IOD)
#define LD_IOD (LD_SET_IOD)
#define LD_ENV (LD_SET_ENV)
#define LD_CANCEL (LD_SET_CANCEL)
// pos 3
#define LD_MEM (LD_SET_MEM)
// implement
#if OMPTARGET_NVPTX_DEBUG
#define DON(_flag) ((unsigned)(OMPTARGET_NVPTX_DEBUG) & (_flag))
#define PRINT0(_flag, _str) \
{ \
if (omptarget_device_environment.debug_level && DON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d>: " _str, blockIdx.x, threadIdx.x, \
threadIdx.x / WARPSIZE, threadIdx.x & 0x1F); \
} \
}
#define PRINT(_flag, _str, _args...) \
{ \
if (omptarget_device_environment.debug_level && DON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d>: " _str, blockIdx.x, threadIdx.x, \
threadIdx.x / WARPSIZE, threadIdx.x & 0x1F, _args); \
} \
}
#else
#define DON(_flag) (FALSE)
#define PRINT0(flag, str)
#define PRINT(flag, str, _args...)
#endif
// for printing without worring about precision, pointers...
#define P64(_x) ((unsigned long long)(_x))
////////////////////////////////////////////////////////////////////////////////
// early defs for test
////////////////////////////////////////////////////////////////////////////////
#define LT_SAFETY (LT_SET_SAFETY | LT_SET_INPUT | LT_SET_FUSSY)
#define LT_INPUT (LT_SET_INPUT | LT_SET_FUSSY)
#define LT_FUSSY (LT_SET_FUSSY)
#if OMPTARGET_NVPTX_TEST == LT_SET_SAFETY
#define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag))
#define ASSERT0(_flag, _cond, _str) \
{ \
if (TON(_flag)) { \
assert(_cond); \
} \
}
#define ASSERT(_flag, _cond, _str, _args...) \
{ \
if (TON(_flag)) { \
assert(_cond); \
} \
}
#elif OMPTARGET_NVPTX_TEST >= LT_SET_INPUT
#define TON(_flag) ((OMPTARGET_NVPTX_TEST) & (_flag))
#define ASSERT0(_flag, _cond, _str) \
{ \
if (TON(_flag) && !(_cond)) { \
printf("<b %3d, t %4d, w %2d, l %2d> ASSERT: " _str "\n", blockIdx.x, \
threadIdx.x, threadIdx.x / WARPSIZE, threadIdx.x & 0x1F); \
assert(_cond); \
} \
}
#define ASSERT(_flag, _cond, _str, _args...) \
{ \
if (TON(_flag) && !(_cond)) { \
printf("<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n", blockIdx.x, \
threadIdx.x, threadIdx.x / WARPSIZE, threadIdx.x & 0x1F, _args); \
assert(_cond); \
} \
}
#else
#define TON(_flag) (FALSE)
#define ASSERT0(_flag, _cond, _str)
#define ASSERT(_flag, _cond, _str, _args...)
#endif
////////////////////////////////////////////////////////////////////////////////
// early defs for warning
#define LW_ALL (LW_SET_ALL)
#define LW_ENV (LW_SET_FUSSY | LW_SET_INPUT | LW_SET_ENV)
#define LW_INPUT (LW_SET_FUSSY | LW_SET_INPUT)
#define LW_FUSSY (LW_SET_FUSSY)
#if OMPTARGET_NVPTX_WARNING
#define WON(_flag) ((OMPTARGET_NVPTX_WARNING) & (_flag))
#define WARNING0(_flag, _str) \
{ \
if (WON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, blockIdx.x, \
threadIdx.x, threadIdx.x / WARPSIZE, threadIdx.x & 0x1F); \
} \
}
#define WARNING(_flag, _str, _args...) \
{ \
if (WON(_flag)) { \
printf("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, blockIdx.x, \
threadIdx.x, threadIdx.x / WARPSIZE, threadIdx.x & 0x1F, _args); \
} \
}
#else
#define WON(_flag) (FALSE)
#define WARNING0(_flag, _str)
#define WARNING(_flag, _str, _args...)
#endif
#endif
|
