Merge branch 'develop' into bf16_verify_tests

.github/workflows/ci.yaml

@@ -351,12 +351,14 @@ jobs:
         docker-images: true
 
     - uses: actions/[email protected]
+      with:
+        fetch-depth: 0 # Fetch the entire repository history and all branches
     - name: Set up Python
       uses: actions/setup-python@v4
       with:
         python-version: 3.8
-    - name: run License Check
-      run: python3 tools/check_stamped.py ${GITHUB_HEAD_REF:-${GITHUB_REF#refs/heads/}}
+    - name: Run License Check
+      run: python3 tools/check_stamped.py origin/${{ github.event_name == 'pull_request' && github.base_ref || 'develop' }}
 
   linux:
 

.github/workflows/config.md

@@ -1,7 +1,7 @@
 #=====ROCM INFO=====
-ROCM_VERSION : '6.0.2'
+ROCM_VERSION : '6.3.1'
 #default ROCm version to be used
-ROCM_BASE_IMAGE : 'rocm/dev-ubuntu-20.04'
+ROCM_BASE_IMAGE : 'rocm/dev-ubuntu-22.04'
 #base image from dockerhub to be used
 ROCM_BUILT_IMAGE : 'rocm-migraphx'
 #name of the docker image built upon ROCm base
@@ -26,4 +26,4 @@ PERFORMANCE_TEST_TIMEOUT : '30m'
 
 #===== W A R N I N G =====
 #VARIABLE NAMES NOT TO BE CHANGED, VALUES ONLY!
-#VALUES MUST BE ENGLOSED IN SINGLE QUOTES!
+#VALUES MUST BE ENGLOSED IN SINGLE QUOTES!

.github/workflows/performance.yaml

@@ -12,7 +12,7 @@ on:
       rocm_release:
         description: ROCm Version
         required: true
-        default: '6.0.2'
+        default: '6.3.1'
       performance_reports_repo:
         description: Repository where performance reports are stored
         required: true
@@ -96,4 +96,4 @@ jobs:
     secrets:
       gh_token: ${{ secrets.MIGRAPHX_BOT_TOKEN }}
       mail_user: ${{ secrets.MAIL_USERNAME }}
-      mail_pass: ${{ secrets.MAIL_PASSWORD }}
+      mail_pass: ${{ secrets.MAIL_PASSWORD }}

CMakeLists.txt

@@ -208,7 +208,6 @@ rocm_enable_clang_tidy(
         -clang-diagnostic-disabled-macro-expansion
         -clang-diagnostic-extern-c-compat
         -clang-diagnostic-unused-command-line-argument
-        -cppcoreguidelines-avoid-capture-default-when-capturing-this
         -cppcoreguidelines-avoid-const-or-ref-data-members
         -cppcoreguidelines-avoid-do-while
         -cppcoreguidelines-explicit-virtual-functions
@@ -222,7 +221,6 @@ rocm_enable_clang_tidy(
         -cppcoreguidelines-pro-type-reinterpret-cast
         -cppcoreguidelines-pro-type-union-access
         -cppcoreguidelines-pro-type-vararg
-        -cppcoreguidelines-rvalue-reference-param-not-moved
         -cppcoreguidelines-special-member-functions
         -cppcoreguidelines-use-default-member-init
         -cppcoreguidelines-virtual-class-destructor

LICENSE

@@ -1,6 +1,6 @@
 The MIT License (MIT)
 
-Copyright (c) 2015-2024 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

codecov.yml

@@ -2,3 +2,4 @@ ignore:
   - "test/"
   - "src/driver"
   - "build/" 
+  - "src/netron_output.cpp"

docs/driver/read.rst

@@ -78,6 +78,10 @@ Print out the program as cpp program.
 
 Print out program as json.
 
+.. option::  --netron
+
+Print out program as a Netron viewable json file.
+
 .. option::  --text
 
 Print out program in text format.

docs/migraphx-driver.rst

@@ -85,6 +85,8 @@ To learn which options can be used with which commands, see the :ref:`MIGraphX d
       - Prints the program in .txt format
    *  - --binary
       - Prints the program in binary format
+   *  - --netron
+      - Prints the program in Netron viewable JSON format
    *  - --output | -o
       - Writes output in a file
    *  - --fill0

docs/reference/driver-options.rst

@@ -16,7 +16,7 @@ read
 
 Loads and prints input graph.
 
-.. include:: ./driver/read.rst
+.. include:: ../driver/read.rst
 
 compile
 -------
@@ -25,8 +25,8 @@ compile
 
 Compiles and prints input graph.
 
-.. include:: ./driver/read.rst
-.. include:: ./driver/compile.rst
+.. include:: ../driver/read.rst
+.. include:: ../driver/compile.rst
 
 run
 ---
@@ -35,8 +35,8 @@ run
 
 Loads and prints input graph.
 
-.. include:: ./driver/read.rst
-.. include:: ./driver/compile.rst
+.. include:: ../driver/read.rst
+.. include:: ../driver/compile.rst
 
 perf
 ----
@@ -45,8 +45,8 @@ perf
 
 Compiles and runs input graph then prints performance report.
 
-.. include:: ./driver/read.rst
-.. include:: ./driver/compile.rst
+.. include:: ../driver/read.rst
+.. include:: ../driver/compile.rst
 
 .. option::  --iterations, -n [unsigned int]
 
@@ -59,8 +59,8 @@ verify
 
 Runs reference and CPU or GPU implementations and checks outputs for consistency.
 
-.. include:: ./driver/read.rst
-.. include:: ./driver/compile.rst
+.. include:: ../driver/read.rst
+.. include:: ../driver/compile.rst
 
 .. option::  --rms-tol [double]
 
@@ -104,5 +104,5 @@ Here is how you can use ``roctx`` combined with :doc:`rocprof <rocprofiler:rocpr
 Running :doc:`rocprof <rocprofiler:rocprofv1>` generates trace information for HIP, HCC and ROCTX in separate ``.txt`` files.
 To understand the interactions between API calls, utilize the :ref:`roctx.py <tools>` helper script.
 
-.. include:: ./driver/read.rst
-.. include:: ./driver/compile.rst
+.. include:: ../driver/read.rst
+.. include:: ../driver/compile.rst

docs/sphinx/requirements.in

@@ -1,2 +1,2 @@
-rocm-docs-core==1.11.0
+rocm-docs-core==1.12.0
 sphinx-collapse

docs/sphinx/requirements.txt

@@ -116,7 +116,7 @@ requests==2.32.3
     # via
     #   pygithub
     #   sphinx
-rocm-docs-core==1.11.0
+rocm-docs-core==1.12.0
     # via -r requirements.in
 smmap==5.0.1
     # via gitdb

examples/diffusion/python_stable_diffusion_3/txt2img.py

@@ -485,6 +485,10 @@ def encode_token_weights(self, model_name, token_weight_pairs):
     def get_embeddings(self, prompt_tokens):
         l_out, l_pooled = self.encode_token_weights("clip-l",
                                                     prompt_tokens["l"])
+        # stable-diffusion-3-lite-onnx has swapped outputs for clip-l text encoder
+        if l_out.shape != (1, 77, 768):
+            l_out, l_pooled = l_pooled, l_out
+
         g_out, g_pooled = self.encode_token_weights("clip-g",
                                                     prompt_tokens["g"])
         if not self.skip_t5:

requirements.txt

@@ -28,4 +28,4 @@ pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build
 msgpack/[email protected] -DMSGPACK_BUILD_TESTS=Off
 sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 ROCm/composable_kernel@b7775add2d28251674d81e220cd4a857b90b997a -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
-ROCm/rocMLIR@13065c4b3a216e1b13dfb8f746b8a0d421f124e8 -DBUILD_FAT_LIBROCKCOMPILER=On
+ROCm/rocMLIR@13065c4b3a216e1b13dfb8f746b8a0d421f124e8 -DBUILD_FAT_LIBROCKCOMPILER=On

src/CMakeLists.txt

@@ -37,6 +37,7 @@ add_library(migraphx
     argument.cpp
     autocast_fp8.cpp
     auto_contiguous.cpp
+    base64.cpp
     common.cpp
     common_dims.cpp
     compile_src.cpp
@@ -73,6 +74,7 @@ add_library(migraphx
     memory_coloring.cpp
     module.cpp
     msgpack.cpp
+    netron_output.cpp
     normalize_attributes.cpp
     normalize_ops.cpp
     op_enums.cpp

src/base64.cpp

@@ -0,0 +1,81 @@
+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2024 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+#include <migraphx/base64.hpp>
+#include <vector>
+#include <array>
+#include <iostream>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+namespace {
+using byte = unsigned char;
+
+std::array<char, 64> constexpr b64_chars{
+    'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P',
+    'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f',
+    'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v',
+    'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/'};
+
+/// base64 encoder snippet altered from https://stackoverflow.com/a/37109258
+std::string encode(const std::vector<byte>& buf)
+{
+    std::size_t len = buf.size();
+    std::vector<byte> res_vec((len + 2) / 3 * 4, '=');
+    std::size_t j         = 0;
+    std::size_t remaining = len % 3;
+    const size_t last     = len - remaining;
+
+    for(size_t i = 0; i < last; i += 3)
+    {
+        std::size_t n = static_cast<std::size_t>(buf.at(i)) << 16u |
+                        static_cast<std::size_t>(buf.at(i + 1)) << 8u |
+                        static_cast<std::size_t>(buf.at(i + 2));
+        res_vec.at(j++) = b64_chars.at(n >> 18u);
+        res_vec.at(j++) = b64_chars.at(n >> 12u & 0x3Fu);
+        res_vec.at(j++) = b64_chars.at(n >> 6u & 0x3Fu);
+        res_vec.at(j++) = b64_chars.at(n & 0x3Fu);
+    }
+    // Set padding
+    if(remaining != 0)
+    {
+        std::size_t n   = --remaining == 0 ? static_cast<std::size_t>(buf.at(last))
+                                           : static_cast<std::size_t>(buf.at(last)) << 8u |
+                                               static_cast<std::size_t>(buf.at(last + 1));
+        res_vec.at(j++) = b64_chars.at(remaining == 0 ? n >> 2u : n >> 10u & 0x3Fu);
+        res_vec.at(j++) = b64_chars.at(remaining == 0 ? n << 4u & 0x3Fu : n >> 4u & 0x03Fu);
+        res_vec.at(j++) = remaining == 0 ? '=' : b64_chars.at(n << 2u & 0x3Fu);
+    }
+    return {res_vec.begin(), res_vec.end()};
+}
+
+} // namespace
+
+std::string base64_encode(const std::string& str)
+{
+    return encode(std::vector<byte>(str.begin(), str.end()));
+}
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/driver/main.cpp

@@ -56,6 +56,8 @@
 #include <migraphx/simplify_reshapes.hpp>
 #include <migraphx/register_target.hpp>
 
+#include <migraphx/netron_output.hpp>
+
 #include <fstream>
 
 namespace migraphx {
@@ -166,6 +168,10 @@ struct loader
            {"--binary"},
            ap.help("Print out program in binary format."),
            ap.set_value("binary"));
+        ap(output_type,
+           {"--netron"},
+           ap.help("Print out program as Netron readable json."),
+           ap.set_value("netron"));
         ap(output, {"--output", "-o"}, ap.help("Output to file."));
     }
 
@@ -418,6 +424,8 @@ struct loader
             *os << to_json_string(p.to_value()) << std::endl;
         else if(type == "binary")
             write(*os, save_buffer(p));
+        else if(type == "netron")
+            *os << make_netron_output(p) << std::endl;
     }
 };
 

src/include/migraphx/base64.hpp

@@ -0,0 +1,39 @@
+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2024 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+#ifndef MIGRAPHX_GUARD_RTGLIB_BASE64_HPP
+#define MIGRAPHX_GUARD_RTGLIB_BASE64_HPP
+
+#include <string>
+#include <migraphx/config.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+/// encode string to base64
+std::string MIGRAPHX_EXPORT base64_encode(const std::string& str);
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/netron_output.hpp

@@ -0,0 +1,39 @@
+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2024 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+#ifndef MIGRAPHX_GUARD_RTGLIB_NETRON_OUTPUT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_NETRON_OUTPUT_HPP
+
+#include <string>
+#include <migraphx/config.hpp>
+#include <migraphx/program.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+MIGRAPHX_EXPORT std::string make_netron_output(const program& prog);
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif