[7.1] bulk merge (#4358)

* Update build and test script and dockerfile to add in onnxrt pai laun… (#4336)

Update build and test script and dockerfile to add in onnxrt pai launcher scripts
Push pai launcher scripts to Onnxrutime to let us reuse scripts from prior run
Related to change from - #4321

* Bump SQlite3 to 3.50.4 (#4322)

* TopK exception bugfix (#4329)

* Use --input-dim for specifying dynamic shapes at driver runtime (#4342)

* Add video and render groups in docker for CI (#4340)

* Fix MXFP4 bugs (#4324)

* propagate_constant ignores `unpack_fp4` instructions now
* `match_find_mx_quantizable_ops` from `simplify_qdq` updated to not require non-constant scales. The scales can be literals.
* transpose, reshape, and broadcast instructions propagated on scale instructions when going to `quant_dot` and `quant_conv`
* `raw_data` `operator<<` updated to use `fallback_visit` to also handle non-computable types

* Brevitas MXFP4 quantization parse (#4301)

* Bump onnx to 1.18.0 (#4323)

* Fuse GEMM+GEMM with rocMLIR (#4261)

rocmlir supports GEMM+GEMM fusion, so MIGraphX needs to support the fusion on our side.
Solves #4230

* Updated SD3 example for change in optimum-onnx[onnxruntime] (#4344)

* Lower lrn to pooling (#4294)

This PR implements a lowering transformation that converts LRN operations into a series of pooling and arithmetic operations.

* Update onnxruntime main be655f69f6c8623eebcac094626d0c8545951e6d (#4348)

* Update docs/reference/MIGraphX-dev-env-vars.rst

---------

Co-authored-by: Ted Themistokleous <[email protected]>
Co-authored-by: Lakhinder Walia <[email protected]>
Co-authored-by: kahmed10 <[email protected]>
Co-authored-by: Charlie Lin <[email protected]>
Co-authored-by: Breanna Devore-McDonald <[email protected]>
Co-authored-by: Aarushi Jain <[email protected]>
Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>

Author: 8 people

Jenkinsfile

@@ -39,6 +39,7 @@ def rocmtestnode(Map conf) {
             export MIGRAPHX_GPU_DEBUG=${gpu_debug}
             export CXX=${compiler}
             export CXXFLAGS='-Werror'
+            rocminfo
             env
             rm -rf build
             mkdir build
@@ -66,12 +67,18 @@ def rocmtestnode(Map conf) {
                 checkout scm
             }
 
+            def video_id = sh(returnStdout: true, script: 'getent group video | cut -d: -f3').trim()
+            def render_id = sh(returnStdout: true, script: 'getent group render | cut -d: -f3').trim()
+            def docker_opts = "--device=/dev/kfd --device=/dev/dri --cap-add SYS_PTRACE -v=${env.WORKSPACE}/../:/workspaces:rw,z"
+            docker_opts = docker_opts + " --group-add=${video_id} --group-add=${render_id} "
+            echo "Docker flags: ${docker_opts}"
+
             gitStatusWrapper(credentialsId: "${env.migraphx_ci_creds}", gitHubContext: "Jenkins - ${variant}", account: 'ROCmSoftwarePlatform', repo: 'AMDMIGraphX') {
                 withCredentials([usernamePassword(credentialsId: 'docker_test_cred', passwordVariable: 'DOCKERHUB_PASS', usernameVariable: 'DOCKERHUB_USER')]) {
                     sh "echo $DOCKERHUB_PASS | docker login --username $DOCKERHUB_USER --password-stdin"
                     pre()
                     sh "docker pull ${DOCKER_IMAGE}:${env.IMAGE_TAG}"
-                    withDockerContainer(image: "${DOCKER_IMAGE}:${env.IMAGE_TAG}", args: "--device=/dev/kfd --device=/dev/dri --group-add video --cap-add SYS_PTRACE -v=${env.WORKSPACE}/../:/workspaces:rw,z ${docker_args}") {
+                    withDockerContainer(image: "${DOCKER_IMAGE}:${env.IMAGE_TAG}", args: docker_opts + docker_args) {
                         timeout(time: 4, unit: 'HOURS') {
                             body(cmake_build)
                         }
@@ -192,7 +199,7 @@ rocmtest clang_debug: rocmnode('mi200+') { cmake_build ->
     }
 }, mlir_debug: rocmnode('mi100+') { cmake_build ->
     stage('MLIR Debug') {
-        withEnv(['MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_MLIR_USE_SPECIFIC_OPS=fused,attention,convolution,dot,convolution_backwards', 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1','MIGRAPHX_DISABLE_LAYERNORM_FUSION=1']) {
+        withEnv(['MIGRAPHX_ENABLE_EXTRA_MLIR=1', 'MIGRAPHX_MLIR_USE_SPECIFIC_OPS=fused,attention,convolution,dot,convolution_backwards', 'MIGRAPHX_ENABLE_MLIR_INPUT_FUSION=1', 'MIGRAPHX_MLIR_ENABLE_SPLITK=1', 'MIGRAPHX_ENABLE_MLIR_REDUCE_FUSION=1', 'MIGRAPHX_ENABLE_MLIR_GEG_FUSION=1', 'MIGRAPHX_ENABLE_SPLIT_REDUCE=1','MIGRAPHX_DISABLE_LAYERNORM_FUSION=1']) {
             def sanitizers = "undefined"
             // Note: the -fno-sanitize= is copied from upstream LLVM_UBSAN_FLAGS.
             def debug_flags = "-g -O2 -fsanitize=${sanitizers} -fno-sanitize=vptr,function -fno-sanitize-recover=${sanitizers}"

dev-requirements.txt

@@ -21,7 +21,7 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 #####################################################################################
-ROCmSoftwarePlatform/rocm-recipes
+ROCm/rocm-recipes
 facebook/[email protected] -X subdir -DCMAKE_DIR=build/cmake
 [email protected] -DENABLE_TESTING=OFF
 pcre,pfultz2/[email protected] -H sha256:d6f7182602a775a7d500a0cedca6449af0400c6493951513046d17615ed0bf11

docs/reference/MIGraphX-dev-env-vars.rst

@@ -144,6 +144,14 @@ Model performance tunable variables change the compilation behavior of a model.
 
       | Default: Reduction fusions are turned off.
 
+  * - | ``MIGRAPHX_ENABLE_MLIR_GEG_FUSION``
+      | Turns on GEMM+GEMM fusions in MLIR.
+    
+    - | ``1``: Turns on G+G fusions.
+      | ``0``: Returns to default behavior.
+
+      | Default: GEMM+GEMM fusions are turned off.
+
   * - | ``MIGRAPHX_MLIR_ENABLE_SPLITK``
       | Turns on Split-k performance configurations during MLIR tuning.
       
@@ -213,6 +221,14 @@ Model performance tunable variables change the compilation behavior of a model.
 
       | Default: No tuning is done for composable kernels.
 
+  * - | ``MIGRAPHX_REWRITE_LRN``
+      | Turns on LRN-to-pooling lowering in the ``rewrite_pooling`` pass.
+
+      
+    - | ``1``: Turns on LRN-to-pooling lowering.
+      | ``0``: Returns to default behavior.
+
+      | Default: LRN-to-pooling lowering is turned off.
                
 Matching
 **********

examples/diffusion/python_stable_diffusion_3/README.md

@@ -17,6 +17,7 @@ python3 -m venv sd_venv
 Install dependencies
 
 ```bash
+pip install --upgrade pip
 pip install -r torch_requirements.txt
 pip install -r requirements.txt
 ```
@@ -37,7 +38,7 @@ huggingface-cli login
 Export the models to onnx. 
 Currently, optimum does not have the changes required in their latest release. Please install from their development branch instead.
 ```bash
-python -m pip install optimum[onnxruntime]@git+https://github.com/huggingface/optimum.git
+pip install "optimum-onnx[onnxruntime]"@git+https://github.com/huggingface/optimum-onnx.git
 ```
 
 Once optimum is built, use the following command to export the models:

rbuild.ini

@@ -2,13 +2,13 @@
 cxx = ${rocm_path}/llvm/bin/clang++
 cc = ${rocm_path}/llvm/bin/clang
 deps =
-    ROCmSoftwarePlatform/rocm-recipes
+    ROCm/rocm-recipes
     -f requirements.txt
 
 [gh]
 ignore =
     danmar/cppcheck
-    ROCmSoftwarePlatform/rocMLIR
+    ROCm/rocMLIR
 deps =
     -f dev-requirements.txt
     oneapi-src/[email protected]

requirements.txt

@@ -27,6 +27,6 @@ nlohmann/[email protected] -DCMAKE_POLICY_VERSION_MINIMUM=3.5
 ROCm/[email protected]
 pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build
 msgpack/[email protected] -DMSGPACK_BUILD_TESTS=Off -DMSGPACK_BUILD_EXAMPLES=Off -DCMAKE_POLICY_VERSION_MINIMUM=3.5
-sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On
+sqlite3@3.50.4 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 ROCm/composable_kernel@b7775add2d28251674d81e220cd4a857b90b997a -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
 ROCm/rocMLIR@33b0fc534532f8e8cb7bec2b5f7d20a69be2def5 -DBUILD_FAT_LIBROCKCOMPILER=On -DLLVM_INCLUDE_TESTS=Off

src/driver/main.cpp

@@ -64,6 +64,9 @@
 #include <iomanip>
 
 namespace {
+
+using dims_map = std::unordered_map<std::string, std::vector<std::size_t>>;
+
 std::vector<std::string>
 get_unrecognized_migraphx_envs(const char* envp[],
                                const std::map<std::string, std::string>& used_env)
@@ -213,7 +216,7 @@ struct loader
 
     static auto parse_param_dims(const std::vector<std::string>& param_dims_info)
     {
-        std::unordered_map<std::string, std::vector<std::size_t>> map_input_dims;
+        dims_map map_input_dims;
         std::string name = "";
         for(auto&& x : param_dims_info)
         {
@@ -502,16 +505,24 @@ struct program_params
         return map_load_args;
     }
 
-    auto generate(const program& p, const target& t, bool offload, unsigned batch)
+    auto generate(const program& p,
+                  const target& t,
+                  bool offload,
+                  unsigned batch,
+                  dims_map map_input_dims = {})
     {
         parameter_map m;
         auto param_shapes = p.get_parameter_shapes();
         std::unordered_map<std::string, shape> static_param_shapes;
-        std::transform(
-            param_shapes.cbegin(),
-            param_shapes.cend(),
-            std::inserter(static_param_shapes, static_param_shapes.end()),
-            [&](const auto& x) { return std::make_pair(x.first, x.second.to_static(batch)); });
+        for(auto&& param : param_shapes)
+        {
+            if(contains(map_input_dims, param.first))
+                static_param_shapes[param.first] = {param.second.type(),
+                                                    map_input_dims[param.first]};
+            else
+                static_param_shapes[param.first] = param.second.to_static(batch);
+        }
+
         for(auto&& s : fill0)
             m[s] = fill_argument(static_param_shapes.at(s), 0);
         for(auto&& s : fill1)
@@ -591,7 +602,8 @@ struct compiler
 
     auto params(const program& p)
     {
-        return parameters.generate(p, ct.get_target(), co.offload_copy, l.batch);
+        return parameters.generate(
+            p, ct.get_target(), co.offload_copy, l.batch, loader::parse_param_dims(l.param_dims));
     }
 
     auto host_params(const program& p)
@@ -730,7 +742,8 @@ struct verify : command<verify>
         std::cout << p << std::endl;
 
         auto t = c.ct.get_target();
-        auto m = c.parameters.generate(p, t, true, c.l.batch);
+        auto m =
+            c.parameters.generate(p, t, true, c.l.batch, loader::parse_param_dims(c.l.param_dims));
 
         if(c.to_fp16)
         {

src/include/migraphx/module.hpp

@@ -339,6 +339,7 @@ struct MIGRAPHX_EXPORT module
     ins_dep_map calc_implicit_deps() const;
 
     void repeat_while_changes(std::size_t n, const std::function<void()>& f);
+    void localized_sort(instruction_ref start_ins, instruction_ref end_ins);
 
     MIGRAPHX_EXPORT friend std::ostream& operator<<(std::ostream& os, const module& m);
     MIGRAPHX_EXPORT friend bool operator==(const module& x, const module& y);

src/include/migraphx/raw_data.hpp

@@ -53,15 +53,15 @@ struct raw_data : raw_data_base
     friend Stream& operator<<(Stream& os, const Derived& d)
     {
         if(not d.empty())
-            d.visit([&](auto x) { os << x; },
-                    [&](auto&& xs) {
-                        for(auto&& x : xs)
-                        {
-                            os << "{ ";
-                            os << x;
-                            os << " }, ";
-                        }
-                    });
+            d.fallback_visit([&](auto x) { os << x; },
+                             [&](auto&& xs) {
+                                 for(auto&& x : xs)
+                                 {
+                                     os << "{ ";
+                                     os << x;
+                                     os << " }, ";
+                                 }
+                             });
         return os;
     }
 
@@ -123,9 +123,13 @@ struct raw_data : raw_data_base
         }
         else
         {
-            auto&& buffer    = static_cast<const Derived&>(*this).data();
+            auto* buffer     = static_cast<const Derived&>(*this).data();
             shape view_shape = {shape::uint8_type, {s.bytes()}};
-            v(make_view(view_shape, reinterpret_cast<byte*>(buffer)));
+            using byte_type =
+                std::conditional_t<std::is_const<std::remove_pointer_t<decltype(buffer)>>{},
+                                   const byte*,
+                                   byte*>;
+            v(make_view(view_shape, reinterpret_cast<byte_type>(buffer)));
         }
     }
 

src/include/migraphx/rewrite_pooling.hpp

@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -38,6 +38,7 @@ struct module;
  */
 struct MIGRAPHX_EXPORT rewrite_pooling
 {
+    bool rewrite_lrn = false;
     std::string name() const { return "rewrite_pooling"; }
     void apply(module& m) const;
 };