Merge remote-tracking branch 'origin/develop' into ref_fp8

docs/contributor_guide.rst

@@ -14,3 +14,4 @@ Contributor Guide
    dev/pass
    dev/matchers
    dev/tools
+   dev/env_vars

docs/dev/env_vars.rst

@@ -0,0 +1,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.

src/driver/verify.cpp

@@ -119,6 +119,7 @@ void verify_program(const std::string& name,
     auto target_outs = run_target(p, t, options, quantize, inputs);
 
     std::size_t output_num = ref_outs.size();
+    bool passed            = true;
     for(std::size_t i = 0; i < output_num; ++i)
     {
         if(ref_outs[i].get_shape().type() != target_outs[i].get_shape().type() or
@@ -130,9 +131,11 @@ void verify_program(const std::string& name,
         }
         else
         {
-            verify_args(name, target_outs[i], verify::expected{ref_outs[i]}, tols);
+            passed &= verify_args(name, target_outs[i], verify::expected{ref_outs[i]}, tols);
         }
     }
+    if(passed)
+        std::cout << "MIGraphX verification passed successfully." << std::endl;
 }
 
 void verify_instructions(const program& prog,

src/include/migraphx/module.hpp

@@ -42,6 +42,7 @@
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 
+MIGRAPHX_EXPORT
 const operation& get_operation(instruction_ref ins);
 
 struct module_impl;

src/include/migraphx/op/isinf.hpp

@@ -35,7 +35,7 @@ struct isinf : unary<isinf>
 {
     auto apply() const
     {
-        return [&](auto x) { return std::isinf(x); };
+        return [&](auto x) { return std::isinf(static_cast<double>(x)); };
     }
 
     std::string name() const { return "isinf"; }

src/include/migraphx/op/slice.hpp

@@ -31,6 +31,7 @@
 #include <migraphx/dyn_output.hpp>
 #include <migraphx/op/normalize_attribute.hpp>
 #include <migraphx/normalize_attributes.hpp>
+#include <array>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/verify_args.cpp

@@ -88,7 +88,6 @@ bool verify_args(const std::string& name,
             if(target_nan_idx >= 0)
                 std::cout << "Non finite number found in target at " << target_nan_idx << ": "
                           << target[target_nan_idx] << std::endl;
-            std::cout << "MIGraphX verification passed successfully." << std::endl;
         }
     });
     return passed;

test/onnx/.onnxrt-commit

@@ -1 +1 @@
-b7b8b5b2ce80edb33990c7ae0fedac6ae3c623f4
+4a8203033930da506b356cdaf88b1531d8d8fca3