Environment Variables

Applies to Linux

2024-05-08

9 min read time

For parsing

MIGRAPHX_TRACE_ONNX_PARSER

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints 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_MATCHES_FOR flag to filter out results.

MIGRAPHX_TRACE_MATCHES_FOR

Set to the name of any matcher to print the traces for that matcher only.

MIGRAPHX_VALIDATE_MATCHES

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Validates 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 output statistics (e.g. min, max, mean). “3” prints everything in “1” and all 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. Prints debug statements for the memory_coloring pass.

MIGRAPHX_TRACE_SCHEDULE

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints 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_8BITS_QUANTIZATION_PARAMS

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints the quantization parameters in the main module only.

MIGRAPHX_DISABLE_DNNL_POST_OPS_WORKAROUND

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Disables the DNNL post ops workaround.

MIGRAPHX_DISABLE_MIOPEN_FUSION

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Disables MIOpen fusions.

MIGRAPHX_DISABLE_SCHEDULE_PASS

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Disables the schedule pass.

MIGRAPHX_DISABLE_REDUCE_FUSION

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Disables the fuse_reduce pass.

MIGRAPHX_ENABLE_SPLIT_REDUCE

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Enable split_reduce.

MIGRAPHX_ENABLE_NHWC

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Enables the layout_nhwc pass.

MIGRAPHX_ENABLE_CK

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Enables use of the Composable Kernels library. Use it in conjunction with MIGRAPHX_DISABLE_MLIR=1.

MIGRAPHX_DISABLE_MLIR*

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Disables use of the rocMLIR library.

MIGRAPHX_COPY_LITERALS

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Uses hip_copy_to_gpu with a new literal instruction rather than using hip_copy_literal{}.

MIGRAPHX_DISABLE_LAYERNORM_FUSION

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Disables layrnorm fusion.

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. Prints trace information for the graph compilation process.

MIGRAPHX_TRACE_PASSES

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints the compile pass and the program after the pass.

MIGRAPHX_TIME_PASSES

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Times the compile passes.

GPU kernels JIT compilation debugging

These environment variables are applicable for both hiprtc and hipclang.

MIGRAPHX_TRACE_CMD_EXECUTE

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints commands executed by the MIGraphX process.

MIGRAPHX_TRACE_HIPRTC

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints HIPRTC options and C++ file executed.

MIGRAPHX_DEBUG_SAVE_TEMP_DIR

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prevents deletion of the created temporary directories.

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. Dumps the compiled HIPRTC source files.

MIGRAPHX_GPU_DUMP_ASM

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Dumps 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. Compiles GPU code in parallel with the given number of threads.

MIGRAPHX_TRACE_NARY

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints the nary device functions used.

MIGRAPHX_ENABLE_HIPRTC_WORKAROUNDS

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Enables HIPRTC workarounds for bugs in HIPRTC.

MIGRAPHX_USE_FAST_SOFTMAX

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Uses fast softmax optimization.

MIGRAPHX_ENABLE_NULL_STREAM

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Allows 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 benchmarking trace. Set to “2” to print detailed benchmarking trace.

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 MLIR operations you want to always use regardless of the GPU architecture. Accepts a list of operators separated by commas (e.g. “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. Performs exhaustive tuning for MLIR.

MIGRAPHX_MLIR_TUNE_LIMIT

Set to an integer greater than 1. Limits the number of solutions available to MLIR for tuning.

CK vars

MIGRAPHX_LOG_CK_GEMM

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints composable kernels GEMM traces.

MIGRAPHX_CK_DEBUG

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Mandatorily adds -DMIGRAPHX_CK_CHECK=1 for compiling composable kernel operators.

MIGRAPHX_TUNE_CK

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Performs tuning for composable kernels.

Testing

MIGRAPHX_TRACE_TEST_COMPILE

Set to the target whose compilation you want to trace (e.g. “gpu”, “cpu”). Prints the compile trace for verify tests on the given target. Don’t use this flag in conjunction with MIGRAPHX_TRACE_COMPILE.

MIGRAPHX_TRACE_TEST

Set to “1”, “enable”, “enabled”, “yes”, or “true” to use. Prints the reference and target programs even if the verify tests pass.

MIGRAPHX_DUMP_TEST

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