env_vars.rst 8.36 KB
Newer Older
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
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
Environment Variables
=====================

For parsing
---------------

**MIGRAPHX_TRACE_ONNX_PARSER**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print debugging traces for the onnx parser.
Prints: initializers (if used), ONNX node operators, added MIGraphX instructions

**MIGRAPHX_DISABLE_FP16_INSTANCENORM_CONVERT**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Disables the conversion from fp16 to fp32 for the InstanceNormalization ONNX operator that MIGX does as a workaround for accuracy issues with reduce_mean/variance.
See ``parse_instancenorm.cpp`` for more details.


Matchers
------------

**MIGRAPHX_TRACE_MATCHES**

Set to "1" to print the matcher that matches an instruction and the matched instruction.
Set to "2" and use the ``MIGRAPHX_TRACE_MATHCES_FOR`` flag to filter out results.

**MIGRAPHX_TRACE_MATCHES_FOR**

Set to the name of any matcher and only traces for that matcher will be printed out.

**MIGRAPHX_VALIDATE_MATCHES**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Validate the module after finding the matches (runs ``module.validate()``).

Program Execution 
---------------------

**MIGRAPHX_TRACE_EVAL**

Set to "1", "2", or "3" to use.
"1" prints the instruction run and the time taken.
"2" prints everything in "1" and a snippet of the output argument and some statistics (ex. min, max, mean) of the output.
"3" prints everything in "1" and the full output buffers.


Program Verification
------------------------

**MIGRAPHX_VERIFY_ENABLE_ALLCLOSE**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Uses ``allclose`` with the given ``atol`` and ``rtol`` for verifying ranges with ``driver verify`` or the tests that use ``migraphx/verify.hpp``.


Pass debugging or Pass controls
-----------------------------------

**MIGRAPHX_TRACE_ELIMINATE_CONTIGUOUS**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Debug print the instructions that have input ``contiguous`` instructions removed.

**MIGRAPHX_DISABLE_POINTWISE_FUSION**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Disables the ``fuse_pointwise`` compile pass.

**MIGRAPHX_DEBUG_MEMORY_COLORING**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print debug statements for the ``memory_coloring`` pass.

**MIGRAPHX_TRACE_SCHEDULE**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print debug statements for the ``schedule`` pass.

**MIGRAPHX_TRACE_PROPAGATE_CONSTANT**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Traces instructions replaced with a constant.

**MIGRAPHX_INT8_QUANTIZATION_PARAMS**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print the quantization parameters in only the main module.

**MIGRAPHX_DISABLE_DNNL_POST_OPS_WORKAROUND**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Disable the DNNL post ops workaround.

**MIGRAPHX_DISABLE_MIOPEN_FUSION**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Disable MIOpen fusions.

**MIGRAPHX_DISABLE_SCHEDULE_PASS**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Disable the ``schedule`` pass.

**MIGRAPHX_DISABLE_REDUCE_FUSION**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Disable the ``fuse_reduce`` pass.

**MIGRAPHX_ENABLE_NHWC**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Enable the ``layout_nhwc`` pass.

**MIGRAPHX_ENABLE_CK**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Enable using the Composable Kernels library.
Should be used in conjunction with ``MIGRAPHX_DISABLE_MLIR=1``.

**MIGRAPHX_DISABLE_MLIR** 
Set to "1", "enable", "enabled", "yes", or "true" to use.
Disable using the rocMLIR library.

**MIGRAPHX_ENABLE_EXTRA_MLIR**
Set to "1", "enable", "enabled", "yes", or "true" to use.
Enables additional opportunities to use MLIR that may improve performance.

**MIGRAPHX_COPY_LITERALS**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Use ``hip_copy_to_gpu`` with a new ``literal`` instruction rather than use ``hip_copy_literal{}``.

Compilation traces
----------------------

**MIGRAPHX_TRACE_FINALIZE**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Debug print instructions during the ``module.finalize()`` step.

**MIGRAPHX_TRACE_COMPILE**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print trace information for the graph compilation process.

**MIGRAPHX_TRACE_PASSES**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print the compile pass and the program after the pass.

**MIGRAPHX_TIME_PASSES**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Time the compile passes.


GPU Kernels JIT compilation debugging (applicable for both hiprtc and hipclang)
-----------------------------------------

**MIGRAPHX_TRACE_CMD_EXECUTE**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print commands executed by the MIGraphX ``process``.

**MIGRAPHX_TRACE_HIPRTC**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print HIPRTC options and C++ file executed.

**MIGRAPHX_DEBUG_SAVE_TEMP_DIR**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Make it so the created temporary directories are not deleted.

**MIGRAPHX_GPU_DEBUG**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Internally, this adds the option ``-DMIGRAPHX_DEBUG`` when compiling GPU kernels. It enables assertions and capture of source locations for the errors. 

**MIGRAPHX_GPU_DEBUG_SYM**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Adds the option ``-g`` when compiling HIPRTC.

**MIGRAPHX_GPU_DUMP_SRC**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Dump the HIPRTC source files compiled.

**MIGRAPHX_GPU_DUMP_ASM**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Dump the hip-clang assembly.

**MIGRAPHX_GPU_OPTIMIZE**

Set the optimization mode for GPU compile (``-O`` option).
Defaults to ``-O3``.

**MIGRAPHX_GPU_COMPILE_PARALLEL**

Set to the number of threads to use.
Compile GPU code in parallel with the given number of threads.

**MIGRAPHX_TRACE_NARY**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print the ``nary`` device functions used.

**MIGRAPHX_ENABLE_HIPRTC_WORKAROUNDS**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Enable HIPRTC workarounds for bugs in HIPRTC.

**MIGRAPHX_USE_FAST_SOFTMAX**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Use the fast softmax optimization.

**MIGRAPHX_ENABLE_NULL_STREAM**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Allow using null stream for miopen and hipStream.

**MIGRAPHX_NSTREAMS**

Set to the number of streams to use.
Defaults to 1.

**MIGRAPHX_TRACE_BENCHMARKING**

Set to "1" to print benchmarching trace.
Set to "2" to print benchmarching trace with more detail.

MLIR vars
-------------

**MIGRAPHX_TRACE_MLIR**

Set to "1" to trace MLIR and print any failures.
Set to "2" to additionally print all MLIR operations.

**MIGRAPHX_MLIR_USE_SPECIFIC_OPS**

Set to the name of the operations you want to always use MLIR regardless of GPU architecture.
Accepts a list of operators separated by commas (ex: "fused", "convolution", "dot").

**MIGRAPHX_MLIR_TUNING_DB**

Set to the path of the MLIR tuning database to load.

**MIGRAPHX_MLIR_TUNING_CFG**

Set to the path of the tuning configuration.
Appends to tuning cfg file that could be used with rocMLIR tuning scripts.

**MIGRAPHX_MLIR_TUNE_EXHAUSTIVE**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Do exhaustive tuning for MLIR.


CK vars
-----------

**MIGRAPHX_LOG_CK_GEMM**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Print Composable Kernels GEMM traces.

**MIGRAPHX_CK_DEBUG**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Always add the ``-DMIGRAPHX_CK_CHECK=1`` for compiling Composable Kernels operators.

**MIGRAPHX_TUNE_CK**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Use tuning for Composable Kernels.

Testing 
------------

**MIGRAPHX_TRACE_TEST_COMPILE**

Set to the target that you want to trace the compilation of (ex. "gpu", "cpu").
Prints the compile trace for the given target for the verify tests.
This flag shouldn't be used in conjunction with ``MIGRAPHX_TRACE_COMPILE``.
For the verify tests only use ``MIGRAPHX_TRACE_TEST_COMPILE``.

**MIGRAPHX_TRACE_TEST**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Prints the reference and target programs even if the verify passed successfully.

**MIGRAPHX_DUMP_TEST**

Set to "1", "enable", "enabled", "yes", or "true" to use.
Dumps verify tests to ``.mxr`` files.