Merge branch 'develop' into dependabot/pip/docs/sphinx/rocm-docs-core…

…-1.10.0
ROCm · Dec 4, 2024 · 696afab · 696afab
2 parents 8603dab + 8576973
commit 696afab
Show file tree

Hide file tree

Showing 47 changed files with 935 additions and 816 deletions.
diff --git a/Jenkinsfile b/Jenkinsfile
@@ -42,7 +42,7 @@ def rocmtestnode(Map conf) {
             rm -rf build
             mkdir build
             cd build
-            cmake -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DBUILD_DEV=On -DCMAKE_EXECUTE_PROCESS_COMMAND_ECHO=STDOUT -DMIGRAPHX_DISABLE_VIRTUAL_ENV=ON ${flags} ..
+            cmake -DCTEST_TIMEOUT=3600 -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DBUILD_DEV=On -DCMAKE_EXECUTE_PROCESS_COMMAND_ECHO=STDOUT -DMIGRAPHX_DISABLE_VIRTUAL_ENV=ON ${flags} ..
             git diff
             git diff-index --quiet HEAD || (echo "Git repo is not clean after running cmake." && exit 1)
             make -j\$(nproc) generate VERBOSE=1
@@ -71,7 +71,7 @@ def rocmtestnode(Map conf) {
                     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=/home/jenkins/:/home/jenkins ${docker_args}") {
-                        timeout(time: 2, unit: 'HOURS') {
+                        timeout(time: 3, unit: 'HOURS') {
                             body(cmake_build)
                         }
                     }

diff --git a/docs/dev/data.rst b/docs/dev/data.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX internal data types
+  :keywords: MIGraphX, code base, contribution, developing, data types
+
 Data types
 ==========
 

diff --git a/docs/dev/dev_intro.rst b/docs/dev/dev_intro.rst
@@ -1,3 +1,8 @@
+.. meta::
+  :description: MIGraphX introduction to developing for the code base
+  :keywords: MIGraphX, code base, contribution, developing, introduction, developers
+
+
 Developer Introduction
 ======================
 

diff --git a/docs/dev/env_vars.rst b/docs/dev/env_vars.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX internal environment variables
+  :keywords: MIGraphX, code base, contribution, developing, env vars, environment variables
+
 Environment Variables
 =====================
 

diff --git a/docs/dev/matchers.rst b/docs/dev/matchers.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX internal matchers
+  :keywords: MIGraphX, code base, contribution, developing, matchers
+
 Matchers
 ========
 

diff --git a/docs/dev/onnx_operators.rst b/docs/dev/onnx_operators.rst
@@ -1,22 +1,24 @@
+.. meta::
+  :description: MIGraphX supported ONNX operators
+  :keywords: MIGraphX, code base, contribution, developing, ONNX operators
+
 Supported ONNX Operators
 ========================
 
 MIGraphX supports operators up to Opset 19. Latest information of ONNX
-operators can be found
-`here <https://github.com/onnx/onnx/blob/master/docs/Operators.md>`__
+operators can be found in `the ONNX GitHub repository <https://github.com/onnx/onnx/blob/master/docs/Operators.md>`_.
 
-MIGraphX supports the following ONNX data types: BOOL, UINT8, UINT16,
-UINT32, UINT64, INT8, INT16, INT32, INT64, FLOAT8, FLOAT16, FLOAT32,
-DOUBLE
+MIGraphX supports the following ONNX data types: BOOL, UINT8, UINT16, UINT32, UINT64, INT8, INT16, INT32, INT64, FLOAT8, FLOAT16, FLOAT32, and DOUBLE
 
-   NOTE: FP8 support is only for E4M3FNUZ, see
-   `here <https://onnx.ai/onnx/technical/float8.html>`__
+   .. Note::
+
+      FP8 support is only for E4M3FNUZ, see `Float stored in 8 bits <https://onnx.ai/onnx/technical/float8.html>`_ in the ONNX documentation.
 
 See below for the support matrix of ONNX operators in MIGraphX.
 
-   NOTE: Supported Types are from ONNX specification. An operator might
-   support more datatypes (e.g. integer type for float operator) than
-   listed.
+   .. Note:: 
+
+      The listed supported types are taken from the ONNX specification. An operator might support other additional datatypes.
 
 Operator Support Matrix
 -----------------------

diff --git a/docs/dev/operators.rst b/docs/dev/operators.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX internal operators
+  :keywords: MIGraphX, code base, contribution, developing, operators
+
 Operators
 =========
 

diff --git a/docs/dev/pass.rst b/docs/dev/pass.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX internal passes
+  :keywords: MIGraphX, code base, contribution, developing, passes
+
 Passes
 ======
 

diff --git a/docs/dev/program.rst b/docs/dev/program.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX program
+  :keywords: MIGraphX, code base, contribution, developing, program
+
 Program
 =======
 

diff --git a/docs/dev/quantization.rst b/docs/dev/quantization.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX internal quantization
+  :keywords: MIGraphX, code base, contribution, developing, quantization
+
 Quantization
 ============
 

diff --git a/docs/dev/targets.rst b/docs/dev/targets.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX targets
+  :keywords: MIGraphX, code base, contribution, developing, targets
+
 Targets
 =======
 

diff --git a/docs/dev/tools.rst b/docs/dev/tools.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: MIGraphX tools
+  :keywords: MIGraphX, code base, contribution, developing, tooks, knobs
+
 .. _tools:
 
 Tools

diff --git a/docs/dev/triage-rocmlir.rst b/docs/dev/triage-rocmlir.rst
@@ -1,3 +1,7 @@
+.. meta::
+  :description: Issue Triaging Guide for suspected issues
+  :keywords: MIGraphX, rocMLIR, issues, pipeline, compilation, bug, code base, kernel, contribution, developing
+
 Issue Triaging Guide for suspected rocMLIR issue
 ================================================
 

diff --git a/docs/index.rst b/docs/index.rst
@@ -33,6 +33,7 @@ The MIGraphX public repository is located at `https://github.com/ROCm/AMDMIGraph
     * :ref:`cpp-api-reference`
     * :ref:`python-api-reference`
     * :ref:`migraphx-driver`
+    * :doc:`Supported ONNX Operators <./dev/onnx_operators>`
 
   .. grid-item-card:: Contributing to the MIGraphX code base
 

diff --git a/docs/reference/cpp.rst b/docs/reference/cpp.rst
@@ -1,3 +1,7 @@
+.. meta::
+   :description: MIGraphX C++ API reference
+   :keywords: MIGraphX, ROCm, C++, API, reference, development, developer
+
 .. _cpp-api-reference:
 
 C++ API reference

diff --git a/docs/reference/driver-options.rst b/docs/reference/driver-options.rst
@@ -1,6 +1,6 @@
 .. meta::
-   :description: MIGraphX provides an optimized execution engine for deep learning neural networks
-   :keywords: MIGraphX, ROCm, library, API, tool
+   :description: MIGraphX driver options
+   :keywords: MIGraphX, ROCm, driver, options
 
 .. _driver-options:
 

diff --git a/docs/reference/py.rst b/docs/reference/py.rst
@@ -1,3 +1,7 @@
+.. meta::
+   :description: MIGraphX Python API reference
+   :keywords: MIGraphX, ROCm, Python, API, reference, development, developer
+
 .. py:module:: migraphx
 
 .. _python-api-reference:

diff --git a/docs/sphinx/_toc.yml.in b/docs/sphinx/_toc.yml.in
@@ -16,6 +16,7 @@ subtrees:
       subtrees:
         - entries:
           - file: reference/driver-options
+    - file: dev/onnx_operators
 
   - caption:  Developing for MIGraphX
     entries:

diff --git a/requirements.txt b/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@45d830dbbc15fe84c41d95585f526a50719020eb -DBUILD_FAT_LIBROCKCOMPILER=On
+ROCm/rocMLIR@c443cf85a09f289c147d7b01f93c1e51390ff65f -DBUILD_FAT_LIBROCKCOMPILER=On
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
@@ -145,6 +145,7 @@ register_migraphx_ops(
     as_shape
     atanh
     atan
+    bit_cast
     bitwise_and
     broadcast
     broadcast_for_dot

diff --git a/src/include/migraphx/op/bit_cast.hpp b/src/include/migraphx/op/bit_cast.hpp
@@ -0,0 +1,104 @@
+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2022 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_OPERATORS_BIT_CAST_HPP
+#define MIGRAPHX_GUARD_OPERATORS_BIT_CAST_HPP
+
+#include <migraphx/op/unary.hpp>
+#include <migraphx/config.hpp>
+#include <migraphx/bit_cast.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+/**
+ * Obtain a value of type `target_type` by reinterpreting
+ * the object represnetaion of the input. Originally used
+ * for casting from fp8e4m3fn to fp8e4m3fnuz.
+ */
+struct bit_cast : unary<bit_cast>
+{
+    shape::type_t target_type;
+
+    template <class Self, class F>
+    static auto reflect(Self& self, F f)
+    {
+        return pack(f(self.target_type, "target_type"));
+    }
+
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs, *this, true}.has(1);
+        auto input = inputs.at(0);
+        std::size_t target_type_size;
+        shape::visit(target_type, [&](auto as) { target_type_size = as.size(); });
+        if(input.type_size() != target_type_size)
+        {
+            MIGRAPHX_THROW("BIT_CAST: target_type has different type_size from input's");
+        }
+        if(input.dynamic())
+        {
+            return {target_type, input.dyn_dims()};
+        }
+        else
+        {
+            return {target_type, input.lens(), input.strides()};
+        }
+    }
+
+    std::string point_op() const
+    {
+        return "${function:bit_cast}<" + shape::cpp_type(target_type) + ">(${0})";
+    }
+
+    argument compute(const dyn_output& dyn_out, std::vector<argument> args) const
+    {
+        argument result{dyn_out.computed_shape};
+        result.visit([&](auto output) {
+            using otype = typename decltype(output)::value_type;
+            args[0].visit([&](auto input) {
+                using itype = typename decltype(input)::value_type;
+                if constexpr(sizeof(otype) == sizeof(itype))
+                {
+                    par_transform(input.begin(), input.end(), output.begin(), [&](auto x) {
+                        return migraphx::bit_cast<otype>(x);
+                    });
+                }
+                else
+                {
+                    // not possible to hit this unless somehow the types change after compute_shape
+                    // is called
+                    MIGRAPHX_THROW("BIT_CAST: type size mismatch");
+                }
+            });
+        });
+        return result;
+    }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif
diff --git a/src/include/migraphx/shape_transform_descriptor.hpp b/src/include/migraphx/shape_transform_descriptor.hpp
@@ -98,7 +98,12 @@ struct MIGRAPHX_EXPORT shape_transform_descriptor
             // the axis. However, it still needs to accounted for. After we
             // generate the broadcast we will set the axis to the hidden
             // axis, and then length to 1.
-            optional<std::size_t> hidden_axis = nullopt;
+            std::vector<std::size_t> hidden_axis = {};
+
+            const std::vector<std::size_t>& origin_axis() const;
+            bool has_hidden_axis() const;
+
+            void add_split_axis(std::size_t i);
 
             MIGRAPHX_EXPORT friend bool operator==(const sub& x, const sub& y);
             MIGRAPHX_EXPORT friend bool operator!=(const sub& x, const sub& y);