Add mlir fusion (#1251)

* Add mlir c api

* Formatting

* Create a type attribute

* Formatting

* Parse module

* Formatting

* Add mlir dump function

* Add test case

* Formatting

* Fix tidy issues

* Update mlit version

* Update to newer mlir

* Format

* Move mlir to the gpu and update the test

* Formatting

* Fix bug when appending module

* Format

* Remove old cmake flag

* Update message

* Add return

* Format

* Add mlir_compile

* Format

* Register dialect

* Handle unsinged integers

* Dont provide output for return instruction

* Format

* Add code to insert memrefs

* Format

* Add mlir verification

* Formatting

* Enable pointwise_fusion

* Disable eliminate_data_type

* Set kernal name

* Format

* Fix device name

* Formatting

* Fix output arg

* Format

* Updates

* Upate hash

* Add fuse_mlir pass

* Format

* Add fuse mlir

* Format

* Update mlir

* Sort parameter names

* Format

* Reenable disabled passes

* Remove old mlir conv

* Remove asym default padding

* Add more verbose tracing

* Format

* Fix compilation errors

* Format

* Whitelist operators

* Format

* Add namespace

* Format

* Update triple

* Format

* Use func dialect

* Format

* Use func.return

* Format

* Upgrade mlir version

* Add comment

* Handle symetrical padding

* Format

* Cleanup debug output

* Format

* List failed tests

* Move mlir compile to jit pipeline

* Format

* Update version

* Add source locations

* Format

* Correctly add module

* Format

* Update failed tests

* Fix failures when mlir is disabled

* Format

* Update mlir version

* Check type for fp32

* Format

* Remove failed test

* Update mlir in driver

* Tidy fixes

* Foramt

* Tidy fixes

* Format

* Fix const

* Remove from requirements

* Fix cmake version

* Fix tidy warning

* Use another ifdef

* Fix tidy

* Other tidy fix

* Format

* Update hash

* Add missing license files

* Format

* Format

* Fix fnction name

Dockerfile

@@ -86,7 +86,7 @@ RUN git clone --single-branch --branch ${ONNXRUNTIME_BRANCH} --recursive ${ONNXR
 
 ADD tools/build_and_test_onnxrt.sh /onnxruntime/build_and_test_onnxrt.sh
 
-RUN PATH=/opt/cmake/bin:$PATH cget -p /usr/local install ROCmSoftwarePlatform/llvm-project-mlir@02078ce236ad90e3aec04c0c770ef5bfc99e49c2
+RUN cget -p /usr/local install ROCmSoftwarePlatform/llvm-project-mlir@26a4b3cfc0a1a15181490f24ae461608fef1b04e -DBUILD_MIXR_TARGET=On
 
 ENV MIOPEN_FIND_DB_PATH=/tmp/miopen/find-db
 ENV MIOPEN_USER_DB_PATH=/tmp/miopen/user-db

src/include/migraphx/matcher.hpp

@@ -349,25 +349,27 @@ match::matcher_result find_match(module& modl, M&& m)
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_TRACE_MATCHES)
 
 /// Find matches for an instruction in the module
-template <class... Ms>
-void find_matches(module& mod, instruction_ref ins, Ms&&... ms)
+template <class Mod, class... Ms>
+void find_matches(Mod& mod, instruction_ref ins, Ms&&... ms)
 {
 #if !defined(__GNUC__) || defined(__clang__) || __GNUC__ > 5
     const
 #endif
-        bool trace = enabled(MIGRAPHX_TRACE_MATCHES{});
-    bool match     = false;
+        int trace = value_of(MIGRAPHX_TRACE_MATCHES{});
+    bool match    = false;
     each_args(
         [&](auto&& m) {
             if(match)
                 return;
-            auto r = match_instruction(mod, ins, m.matcher());
-            if(r.result == mod.end())
+            if(trace > 1)
+                std::cout << "Match: " << get_type_name(m) << std::endl;
+            auto r = match_instruction(get_module(mod), ins, m.matcher());
+            if(r.result == get_module(mod).end())
                 return;
-            if(trace)
+            if(trace > 0)
             {
                 std::cout << "Matched by " << get_type_name(m) << std::endl;
-                mod.debug_print(ins);
+                get_module(mod).debug_print(ins);
             }
             m.apply(mod, r);
             match = true;
@@ -376,10 +378,10 @@ void find_matches(module& mod, instruction_ref ins, Ms&&... ms)
 }
 
 /// Find matches in a module
-template <class... Ms>
-void find_matches(module& mod, Ms&&... ms)
+template <class Mod, class... Ms>
+void find_matches(Mod& mod, Ms&&... ms)
 {
-    for(auto ins : iterator_for(mod))
+    for(auto ins : iterator_for(get_module(mod)))
     {
         find_matches(mod, ins, ms...);
     }

src/include/migraphx/module.hpp

@@ -124,7 +124,7 @@ struct module
                      std::unordered_map<instruction_ref, instruction_ref> map_ins = {});
 
     std::vector<instruction_ref>
-    add_instructions(module_ref m,
+    add_instructions(const_module_ref m,
                      std::unordered_map<instruction_ref, instruction_ref> map_ins = {});
 
     std::vector<instruction_ref>
@@ -139,7 +139,7 @@ struct module
 
     std::vector<instruction_ref>
     insert_instructions(instruction_ref ins,
-                        module_ref m,
+                        const_module_ref m,
                         std::unordered_map<instruction_ref, instruction_ref> map_ins = {});
 
     std::vector<instruction_ref>
@@ -227,6 +227,8 @@ struct module
     std::unique_ptr<module_impl> impl;
 };
 
+inline module& get_module(module& m) { return m; }
+
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
 

src/include/migraphx/module_ref.hpp

@@ -32,7 +32,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 
 struct module;
-using module_ref = module*;
+using module_ref       = module*;
+using const_module_ref = const module*;
 
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/include/migraphx/ranges.hpp

@@ -198,6 +198,12 @@ void transform(Range&& r, Iterator it, F f)
     std::transform(r.begin(), r.end(), it, f);
 }
 
+template <class Range1, class Range2, class Iterator, class F>
+void transform(Range1&& r1, Range2&& r2, Iterator it, F f)
+{
+    std::transform(r1.begin(), r1.end(), r2.begin(), it, f);
+}
+
 template <class Range>
 auto reverse(Range& r)
 {

src/include/migraphx/shape.hpp

@@ -191,6 +191,10 @@ struct shape
 
         std::size_t size(std::size_t n = 1) const { return sizeof(type) * n; }
 
+        auto is_integral() const { return std::is_integral<type>{}; }
+        auto is_signed() const { return std::is_signed<type>{}; }
+        auto is_unsigned() const { return std::is_unsigned<type>{}; }
+
         template <class U>
         type* from(U* buffer, std::size_t n = 0) const
         {

src/module.cpp

@@ -399,7 +399,8 @@ module::add_instructions(const std::vector<instruction_ref>& instructions,
 }
 
 std::vector<instruction_ref>
-module::add_instructions(module_ref m, std::unordered_map<instruction_ref, instruction_ref> map_ins)
+module::add_instructions(const_module_ref m,
+                         std::unordered_map<instruction_ref, instruction_ref> map_ins)
 {
     return this->insert_instructions(this->end(), m, std::move(map_ins));
 }
@@ -420,8 +421,10 @@ module::insert_instructions(instruction_ref ins,
     return insert_generic_instructions(*this, ins, instructions, std::move(map_ins));
 }
 
-std::vector<instruction_ref> module::insert_instructions(
-    instruction_ref ins, module_ref m, std::unordered_map<instruction_ref, instruction_ref> map_ins)
+std::vector<instruction_ref>
+module::insert_instructions(instruction_ref ins,
+                            const_module_ref m,
+                            std::unordered_map<instruction_ref, instruction_ref> map_ins)
 {
     return insert_generic_instructions(*this, ins, iterator_for(*m), std::move(map_ins));
 }

src/targets/gpu/CMakeLists.txt

@@ -164,6 +164,7 @@ add_library(migraphx_gpu
     deconvolution.cpp
     device_name.cpp
     elu.cpp
+    fuse_mlir.cpp
     fuse_ops.cpp
     gather.cpp
     gemm_impl.cpp
@@ -176,7 +177,7 @@ add_library(migraphx_gpu
     loop.cpp
     lrn.cpp
     leaky_relu.cpp
-    mlir_conv.cpp
+    mlir.cpp
     multinomial.cpp
     nonzero.cpp
     pack_args.cpp
@@ -320,16 +321,26 @@ message(STATUS "extractkernel: ${MIGRAPHX_EXTRACT_KERNEL}")
 
 set(MIGRAPHX_ENABLE_MLIR OFF CACHE BOOL "")
 if(MIGRAPHX_ENABLE_MLIR)
-    find_library(LIBMLIRMIOPEN MLIRMIOpenThin REQUIRED)
+    find_library(MLIRAPI_LIBRARY MLIRMIOpen 
+        PATH_SUFFIXES
+        # Workaournd broken mlir install
+        lib/ lib/lib)
     # REQUIRED is not supported before cmake 3.18
-    if(NOT LIBMLIRMIOPEN)
-        message(FATAL_ERROR "libMLIRMIOpenThin not found")
+    if(NOT MLIRAPI_LIBRARY)
+        message(FATAL_ERROR "libMLIRMIOpen not found")
     else()
-        message(STATUS "Build with libMLIRMIOpenThin: " ${LIBMLIRMIOPEN})
+        message(STATUS "Build with libMLIRMIOpen: " ${MLIRAPI_LIBRARY})
     endif()
 
-    target_compile_definitions(migraphx_gpu PRIVATE "-DMIGRAPHX_MLIR_MIOPEN_SUPPORT")
-    target_link_libraries(migraphx_gpu PUBLIC ${LIBMLIRMIOPEN})
+    find_path(MLIRAPI_HEADERS NAMES mlir-c/Dialect/MIGraphX.h)
+    # Workaround MLIR broken installation
+    find_path(MLIRAPI_HEADERS2 NAMES mlir-c/Registration.h
+        PATH_SUFFIXES 
+        include/external/include external/include)
+
+    target_compile_definitions(migraphx_gpu PRIVATE "-DMIGRAPHX_MLIR")
+    target_include_directories(migraphx_gpu SYSTEM PRIVATE ${MLIRAPI_HEADERS} ${MLIRAPI_HEADERS2})
+    target_link_libraries(migraphx_gpu PUBLIC ${MLIRAPI_LIBRARY})
 endif()
 
 set(MIGRAPHX_USE_HIPRTC OFF CACHE BOOL "")

src/targets/gpu/code_object_op.cpp

@@ -52,7 +52,7 @@ code_object_op::compute(context& ctx, const shape&, const std::vector<argument>&
     std::transform(
         args.begin(), args.end(), kargs.begin(), [](const argument& a) { return a.data(); });
     k.launch(ctx.get_stream().get(), global, local, std::move(kargs));
-    return args.back();
+    return args[get_output_arg(args.size())];
 }
 void code_object_op::finalize(context&, const shape&, const std::vector<shape>&)
 {

src/targets/gpu/fuse_mlir.cpp

+/*
+ * 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.
+ */
+#include <migraphx/gpu/fuse_mlir.hpp>
+#include <migraphx/gpu/mlir.hpp>
+#include <migraphx/matcher.hpp>
+#include <migraphx/pass_manager.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/register_op.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+struct module;
+
+namespace gpu {
+
+#ifdef MIGRAPHX_MLIR
+struct mlir_conv
+{
+    operation op = make_op("convolution");
+
+    template <class Self, class F>
+    static auto reflect(Self& self, F f)
+    {
+        return pack(f(self.op, "op"));
+    }
+
+    std::string name() const { return "gpu::mlir_conv"; }
+    shape compute_shape(std::vector<shape> inputs, const std::vector<module_ref>& mods) const
+    {
+        check_shapes{inputs, *this}.standard();
+        if(mods.size() != 1)
+            MIGRAPHX_THROW("should have one submodule.");
+        if(inputs.size() < 2)
+            MIGRAPHX_THROW("should have at least two inputs.");
+        auto n = inputs.size();
+        return op.compute_shape({inputs[n - 2], inputs[n - 1]});
+    }
+};
+MIGRAPHX_REGISTER_OP(mlir_conv);
+
+namespace {
+struct find_conv_pointwise
+{
+    // Find a convolution followed by a pointwise operation.
+    auto matcher() const
+    {
+        auto convolution =
+            match::skip(match::name("contiguous"))(match::name("convolution").bind("convolution"));
+        return match::name("pointwise")(match::any_of[match::inputs()](convolution.bind("x")));
+    }
+
+    void apply(module_pass_manager& mpm, const match::matcher_result& r) const
+    {
+        auto ins      = r.result;
+        auto conv_ins = r.instructions["convolution"];
+        auto x_ins    = r.instructions["x"]; // input after contiguous
+        auto* pm      = ins->module_inputs().front();
+        auto names    = pm->get_parameter_names();
+        // Whitelist pointwise operators
+        if(std::any_of(pm->begin(), pm->end(), [](const auto& i) {
+               return not contains({"@literal", "@param", "@return", "convolution", "add", "relu"},
+                                   i.name());
+           }))
+            return;
+        // Only fuse with fp32 for now
+        if(std::any_of(ins->inputs().begin(), ins->inputs().end(), [&](auto i) {
+               return i->get_shape().type() != shape::type_t::float_type;
+           }))
+            return;
+        std::sort(names.begin(), names.end());
+        module_ref mm = mpm.create_module("mlir_" + pm->name());
+        mm->set_bypass();
+        std::unordered_map<instruction_ref, instruction_ref> param_map;
+        auto x    = mm->add_parameter("x" + std::to_string(names.size()),
+                                   conv_ins->inputs().at(0)->get_shape());
+        auto w    = mm->add_parameter("x" + std::to_string(names.size() + 1),
+                                   conv_ins->inputs().at(1)->get_shape());
+        auto conv = mm->add_instruction(conv_ins->get_operator(), {x, w});
+        std::transform(names.begin(),
+                       names.end(),
+                       ins->inputs().begin(),
+                       std::inserter(param_map, param_map.end()),
+                       [&](auto name, auto input) {
+                           if(input == x_ins)
+                               return std::make_pair(pm->get_parameter(name), conv);
+                           return std::make_pair(pm->get_parameter(name),
+                                                 mm->add_parameter(name, input->get_shape()));
+                       });
+        mm->add_return(mm->insert_instructions(mm->end(), pm, param_map));
+
+        std::vector<instruction_ref> inputs;
+        std::copy_if(ins->inputs().begin(),
+                     ins->inputs().end(),
+                     std::back_inserter(inputs),
+                     [&](auto input) { return input != conv_ins; });
+        inputs.insert(inputs.end(), conv_ins->inputs().begin(), conv_ins->inputs().end());
+        mpm.get_module().replace_instruction(
+            ins, mlir_conv{conv_ins->get_operator()}, inputs, {mm});
+    }
+};
+} // namespace
+
+#endif
+
+void fuse_mlir::apply(module_pass_manager& mpm) const
+{
+#ifdef MIGRAPHX_MLIR
+    match::find_matches(mpm, find_conv_pointwise{});
+#else
+    (void)mpm;
+#endif
+}
+
+} // namespace gpu
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/fuse_ops.cpp

@@ -336,6 +336,7 @@ void move_standard_front(std::vector<instruction_ref>& args)
 
 auto gpu_name(const std::string& s) { return match::name("gpu::" + s); }
 
+namespace {
 struct find_layernorm
 {
     auto matcher() const { return match::layernorm(&gpu_name); }
@@ -836,15 +837,6 @@ inline auto precompile_name(std::string s) // NOLINT
     });
 }
 
-template <class... Ms>
-auto conv_bias_pointwise(Ms... ms)
-{
-    return precompile_name("pointwise")(
-        match::either_arg(0, 1)(bias_shape(match::used_once()).bind("bias"),
-                                fusable_conv(match::used_once()).bind("conv")),
-        ms...);
-}
-
 struct find_conv_bias
 {
     context* ctx = nullptr;
@@ -1013,6 +1005,7 @@ struct find_commutative_broadcast
         m.replace_instruction(ins, ins->get_operator(), args);
     }
 };
+} // namespace
 
 struct find_contiguous
 {

src/targets/gpu/include/migraphx/gpu/code_object_op.hpp

@@ -38,12 +38,13 @@ struct context;
 
 struct code_object_op
 {
-    value::binary code_object;
-    std::string symbol_name;
-    std::size_t global;
-    std::size_t local;
-    std::vector<shape> expected_inputs;
-    shape output;
+    value::binary code_object{};
+    std::string symbol_name = "";
+    std::size_t global      = 0;
+    std::size_t local       = 0;
+    std::vector<shape> expected_inputs{};
+    shape output{};
+    std::int64_t output_arg = -1;
     kernel k{};
 
     template <class Self, class F>
@@ -66,9 +67,13 @@ struct code_object_op
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
     void finalize(context&, const shape&, const std::vector<shape>&);
+    std::int64_t get_output_arg(std::size_t n) const
+    {
+        return output_arg < 0 ? n + output_arg : output_arg;
+    }
     std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
     {
-        return shapes.size() - 1;
+        return get_output_arg(shapes.size());
     }
 
     friend std::ostream& operator<<(std::ostream& os, const code_object_op& op)

src/targets/gpu/include/migraphx/gpu/mlir_conv.hpp → src/targets/gpu/include/migraphx/gpu/fuse_mlir.hpp

@@ -21,8 +21,8 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
  * THE SOFTWARE.
  */
-#ifndef MIGRAPHX_GUARD_RTGLIB_MIOPEN_MLIR_CONV_HPP
-#define MIGRAPHX_GUARD_RTGLIB_MIOPEN_MLIR_CONV_HPP
+#ifndef MIGRAPHX_GUARD_GPU_FUSE_MLIR_HPP
+#define MIGRAPHX_GUARD_GPU_FUSE_MLIR_HPP
 
 #include <migraphx/config.hpp>
 #include <migraphx/gpu/context.hpp>
@@ -30,18 +30,19 @@
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 
-struct module;
+struct module_pass_manager;
 
 namespace gpu {
-struct mlir_conv
+
+struct fuse_mlir
 {
-    context* ctx;
-    std::string name() const { return "mlir::convolution"; }
-    void apply(module& m) const;
+    context* ctx = nullptr;
+    std::string name() const { return "gpu::fuse_mlir"; }
+    void apply(module_pass_manager& mpm) const;
 };
 
 } // namespace gpu
+
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
-
-#endif
+#endif // MIGRAPHX_GUARD_GPU_FUSE_MLIR_HPP

src/targets/gpu/include/migraphx/gpu/mlir.hpp

+/*
+ * 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_RTGLIB_GPU_MLIR_HPP
+#define MIGRAPHX_GUARD_RTGLIB_GPU_MLIR_HPP
+
+#include <string>
+#include <vector>
+#include <migraphx/config.hpp>
+#include <migraphx/gpu/code_object_op.hpp>
+#include <migraphx/instruction_ref.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+struct module;
+namespace gpu {
+
+std::string dump_mlir(const module& m);
+code_object_op compile_mlir(const context& ctx, const module& m);
+
+instruction_ref insert_mlir(module& m,
+                            instruction_ref ins,
+                            code_object_op co,
+                            const std::vector<instruction_ref>& inputs);
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/jit/mlir.cpp

+/*
+ * 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.
+ */
+#include <migraphx/gpu/compiler.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/gpu/context.hpp>
+
+#include <migraphx/gpu/mlir.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct mlir_compiler : compiler<mlir_compiler>
+{
+    std::vector<std::string> names() const { return {"gpu::mlir_conv"}; }
+
+    operation compile_op(context&, const std::vector<shape>&, const value&) const { return {}; }
+
+    compiler_replace compile(context& ctx, instruction_ref ins, const operation&) const
+    {
+        auto* smod = ins->module_inputs().front();
+        assert(smod->get_parameter_names().size() == ins->inputs().size() - 1);
+        return insert(compile_mlir(ctx, *smod));
+    }
+
+    compiler_replace insert(code_object_op co) const
+    {
+        return [co = std::move(co)](module& m, instruction_ref ins) {
+            auto mlir = insert_mlir(m, ins, co, ins->inputs());
+            m.replace_instruction(ins, mlir);
+        };
+    }
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/mlir_conv.cpp

-/*
- * 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.
- */
-#include <migraphx/gpu/mlir_conv.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/instruction.hpp>
-#include <migraphx/make_op.hpp>
-
-#include <migraphx/op/convolution.hpp>
-
-#include <migraphx/gpu/context.hpp>
-#include <migraphx/gpu/convolution.hpp>
-#include <migraphx/iterator_for.hpp>
-#include <migraphx/program.hpp>
-
-#include <migraphx/ranges.hpp>
-#include <migraphx/make_op.hpp>
-#include <migraphx/generate.hpp>
-#include <migraphx/program.hpp>
-#include <migraphx/gpu/kernel.hpp>
-#include <migraphx/gpu/target.hpp>
-#include <migraphx/gpu/hip.hpp>
-#include <migraphx/gpu/compile_hip.hpp>
-
-#include <utility>
-#include <functional>
-#include <algorithm>
-
-#ifdef MIGRAPHX_MLIR_MIOPEN_SUPPORT
-#include <Miir.h>
-#endif // MIGRAPHX_MLIR_MIOPEN_SUPPORT
-
-#include <cstdio>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-struct mlir_apply
-{
-    module* mod           = nullptr;
-    const mlir_conv* pass = nullptr;
-
-    const char* mlir_kernel_name = "migraphx_conv2d";
-
-    std::unordered_map<uint64_t, instruction_ref> literal_map{};
-
-    struct execution_spec
-    {
-        migraphx::value::binary binary;
-        size_t global_size;
-        size_t local_size;
-        execution_spec(migraphx::value::binary&& binary_m, size_t global_s, size_t local_s)
-            : binary(std::move(binary_m)), global_size(global_s), local_size(local_s)
-        {
-        }
-    };
-
-    std::unordered_map<std::string, std::shared_ptr<execution_spec>> binary_map{};
-
-    context& get_context() const
-    {
-        assert(pass != nullptr);
-        assert(pass->ctx != nullptr);
-        return *pass->ctx;
-    }
-
-    void init() const
-    {
-        assert(mod != nullptr);
-        assert(pass != nullptr);
-    }
-
-    std::shared_ptr<execution_spec> make_mlir_binary(instruction_ref op_r)
-    {
-        std::shared_ptr<execution_spec> result;
-
-#ifdef MIGRAPHX_MLIR_MIOPEN_SUPPORT
-        auto conv  = any_cast<op::convolution>(op_r->get_operator());
-        auto inp_t = op_r->inputs().at(0)->get_shape();
-        auto flt_t = op_r->inputs().at(1)->get_shape();
-        auto out_t = op_r->get_shape();
-
-        auto get_type_str = [](const shape& s) -> const char* {
-            switch(s.type())
-            {
-            case shape::float_type: return "f32";
-            case shape::half_type: return "f16";
-            case shape::bool_type:
-            case shape::double_type:
-            case shape::uint8_type:
-            case shape::int8_type:
-            case shape::uint16_type:
-            case shape::int16_type:
-            case shape::int32_type:
-            case shape::int64_type:
-            case shape::uint32_type:
-            case shape::uint64_type:
-            case shape::tuple_type: break;
-            }
-            return nullptr;
-        };
-
-        const auto* inp_t_s = get_type_str(inp_t);
-        const auto* flt_t_s = get_type_str(flt_t);
-        const auto* out_t_s = get_type_str(out_t);
-
-        if(out_t_s == nullptr || inp_t_s == nullptr || flt_t_s == nullptr)
-            return result;
-
-        std::string mlir_options = "--kernel_name " + std::string(mlir_kernel_name);
-
-        // platform spec
-        auto& device = get_context().get_current_device();
-        char dev_name[64];
-        sprintf(dev_name, "gfx%lu%02lu", device.get_device_major(), device.get_device_minor());
-        mlir_options += " --arch " + std::string(dev_name) + " --num_cu " +
-                        std::to_string(device.get_cu_count()); // ???
-
-        // Conv spec
-        mlir_options +=
-            " --operation "
-            "conv2d"
-            " --batchsize " +
-            std::to_string(conv.group) + " --groupsize " + std::to_string(1) + " --padding_h " +
-            std::to_string(conv.padding[0]) + " --padding_w " + std::to_string(conv.padding[1]) +
-            " --conv_stride_h " + std::to_string(conv.stride[0]) + " --conv_stride_w " +
-            std::to_string(conv.stride[1]) + " --dilation_h " + std::to_string(conv.dilation[0]) +
-            " --dilation_w " + std::to_string(conv.dilation[1]);
-
-        // Input spec
-        mlir_options += " --in_layout "
-                        "NCHWG"
-                        " --in_type " +
-                        std::string(inp_t_s) + " --in_channels " + std::to_string(inp_t.lens()[1]) +
-                        " --in_h " + std::to_string(inp_t.lens()[2]) + " --in_w " +
-                        std::to_string(inp_t.lens()[3]);
-
-        // Filter spec
-        mlir_options += " --fil_layout "
-                        "NCHWG"
-                        " --fil_type " +
-                        std::string(flt_t_s) + " --fil_h " + std::to_string(flt_t.lens()[2]) +
-                        " --fil_w " + std::to_string(flt_t.lens()[3]);
-
-        // Output spec
-        mlir_options += " --out_layout "
-                        "NCHWG"
-                        " --out_type " +
-                        std::string(out_t_s) + " --out_channels " +
-                        std::to_string(out_t.lens()[1]) + " --out_h " +
-                        std::to_string(out_t.lens()[2]) + " --out_w " +
-                        std::to_string(out_t.lens()[3]);
-
-        auto bin_i = binary_map.find(mlir_options);
-        if(bin_i == binary_map.end())
-        {
-            size_t bin_size = 0;
-
-            using mlir_handle = MIGRAPHX_MANAGE_PTR(MiirHandle, miirDestroyHandle);
-            auto handle       = mlir_handle(miirCreateHandle(mlir_options.c_str()));
-
-            if(miirLowerBin(handle.get()) == MIIR_SUCCESS &&
-               miirBufferGet(handle.get(), nullptr, &bin_size) == MIIR_SUCCESS)
-            {
-                migraphx::value::binary bin(bin_size);
-                if(miirBufferGet(handle.get(), reinterpret_cast<char*>(bin.data()), &bin_size) ==
-                   MIIR_SUCCESS)
-                {
-                    size_t global_size;
-                    size_t block_size;
-                    if(miirGetExecutionDims(handle.get(), &global_size, &block_size) ==
-                       MIIR_SUCCESS)
-                    {
-                        result = std::make_shared<execution_spec>(
-                            std::move(bin), global_size, block_size);
-                    }
-                }
-            }
-
-            binary_map[mlir_options] = result;
-        }
-        else
-        {
-            result = bin_i->second;
-        }
-#else  // MIGRAPHX_MLIR_MIOPEN_SUPPORT
-        (void)op_r;
-#endif // MIGRAPHX_MLIR_MIOPEN_SUPPORT
-        return result;
-    }
-
-    instruction_ref get_literal(uint64_t value)
-    {
-        auto fi = literal_map.find(value);
-        if(fi != literal_map.end())
-            return fi->second;
-        auto lit = mod->add_literal(value);
-        literal_map.emplace(value, lit);
-        return lit;
-    }
-
-    operation make_code_object_op(instruction_ref op_r, const std::shared_ptr<execution_spec>& spec)
-    {
-        // each pointer is expanded out to a MemRefDescriptor
-        auto inp_t = op_r->inputs().at(0)->get_shape();
-        auto flt_t = op_r->inputs().at(1)->get_shape();
-        auto out_t = op_r->get_shape();
-
-        auto i64 = shape(shape::uint64_type);
-
-        std::vector<shape> expected_inputs = {
-            flt_t, flt_t, i64, i64, i64, i64, i64, i64, i64, i64, i64, i64,  i64,   inp_t,
-            inp_t, i64,   i64, i64, i64, i64, i64, i64, i64, i64, i64, i64,  out_t, out_t,
-            i64,   i64,   i64, i64, i64, i64, i64, i64, i64, i64, i64, out_t};
-
-        return migraphx::make_op("gpu::code_object",
-                                 {
-                                     {"code_object", spec->binary},
-                                     {"symbol_name", mlir_kernel_name},
-                                     {"global", spec->global_size},
-                                     {"local", spec->local_size},
-                                     {"expected_inputs", migraphx::to_value(expected_inputs)},
-                                     {"output", migraphx::to_value(out_t)},
-                                 });
-    }
-
-    void add_memref_descriptor(std::vector<instruction_ref>& refs, instruction_ref inst)
-    {
-        const size_t offset = 0;
-        auto inst_t         = inst->get_shape();
-        refs.push_back(inst);
-        refs.push_back(inst);
-        refs.push_back(get_literal(offset)); // offset
-
-        // dim sizes
-        std::transform(inst_t.lens().begin(),
-                       inst_t.lens().end(),
-                       std::back_inserter(refs),
-                       [&](const auto& lval) { return get_literal(lval); });
-        refs.push_back(get_literal(1)); // G
-
-        // dim strides
-        std::transform(inst_t.strides().begin(),
-                       inst_t.strides().end(),
-                       std::back_inserter(refs),
-                       [&](const auto& lval) { return get_literal(lval); });
-        refs.push_back(get_literal(1)); // G
-    }
-
-    instruction_ref insert_allocation(instruction_ref ins, const shape& s) const
-    {
-        return mod->insert_instruction(ins, hip_allocate{s});
-    }
-
-    void replace_conv_op(instruction_ref ins)
-    {
-        auto conv_bin = make_mlir_binary(ins);
-        if(conv_bin)
-        {
-            auto conv = make_code_object_op(ins, conv_bin);
-
-            auto inp = ins->inputs().at(0);
-            auto flt = ins->inputs().at(1);
-            auto out = insert_allocation(ins, ins->get_shape());
-
-            std::vector<instruction_ref> refs;
-            refs.reserve(3 * 13 + 1);
-            add_memref_descriptor(refs, flt);
-            add_memref_descriptor(refs, inp);
-            add_memref_descriptor(refs, out);
-            refs.push_back(out);
-
-            mod->replace_instruction(ins, conv, refs);
-        }
-    }
-
-    void apply()
-    {
-        init();
-        for(auto it : iterator_for(*mod))
-        {
-            if(it->name() == "convolution")
-            {
-                replace_conv_op(it);
-            }
-        }
-    }
-};
-
-void mlir_conv::apply(module& m) const { mlir_apply{&m, this}.apply(); }
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

src/targets/gpu/target.cpp

@@ -53,10 +53,10 @@
 #include <migraphx/gpu/compile_ops.hpp>
 #include <migraphx/gpu/concat_gpu_opt.hpp>
 #include <migraphx/gpu/context.hpp>
+#include <migraphx/gpu/fuse_mlir.hpp>
 #include <migraphx/gpu/fuse_ops.hpp>
 #include <migraphx/gpu/prefuse_ops.hpp>
 #include <migraphx/gpu/lowering.hpp>
-#include <migraphx/gpu/mlir_conv.hpp>
 #include <migraphx/gpu/pack_int8_args.hpp>
 #include <migraphx/gpu/schedule_model.hpp>
 #include <migraphx/gpu/sync_device.hpp>
@@ -128,7 +128,8 @@ std::vector<pass> target::get_passes(migraphx::context& gctx, const compile_opti
         dead_code_elimination{},
         enable_pass(not enabled(MIGRAPHX_DISABLE_POINTWISE_FUSION{}), fuse_pointwise{}),
         dead_code_elimination{},
-        mlir_conv{&ctx},
+        fuse_mlir{&ctx},
+        dead_code_elimination{},
         lowering{&ctx, options.offload_copy},
         eliminate_contiguous{"gpu::contiguous"},
         dead_code_elimination{},

test/gpu/mlir.cpp

+/*
+ * 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.
+ */
+#include <migraphx/gpu/mlir.hpp>
+#include <migraphx/gpu/target.hpp>
+#include <migraphx/gpu/context.hpp>
+#include <migraphx/gpu/write_literals.hpp>
+#include <migraphx/ref/target.hpp>
+#include <migraphx/module.hpp>
+#include <migraphx/program.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/verify_args.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/functional.hpp>
+#include <test.hpp>
+
+using migraphx::trim;
+
+// m test_gpu_mlir && ./bin/test_gpu_mlir
+
+struct mlir_gpu_target : migraphx::gpu::target
+{
+    std::string name() const { return "mlir"; }
+    std::vector<migraphx::pass> get_passes(migraphx::context& gctx,
+                                           const migraphx::compile_options&) const
+    {
+        auto& ctx = migraphx::any_cast<migraphx::gpu::context>(gctx);
+        return {migraphx::gpu::write_literals{&ctx}};
+    }
+};
+
+std::string encode(const std::string& s)
+{
+    std::stringstream ss;
+    bool prespace = false;
+    for(auto c : s)
+    {
+        if(std::isspace(c) != 0)
+        {
+            if(not prespace)
+                ss << "  ";
+            prespace = true;
+        }
+        else if(std::isprint(c) != 0)
+        {
+            ss << c;
+            prespace = false;
+        }
+    }
+    return migraphx::trim(ss.str());
+}
+
+migraphx::program create_program_from_mlir(const migraphx::module& mmlir)
+{
+    migraphx::program p;
+    auto* mm   = p.get_main_module();
+    auto names = mmlir.get_parameter_names();
+    std::vector<migraphx::instruction_ref> inputs;
+    std::transform(names.begin(), names.end(), std::back_inserter(inputs), [&](const auto& name) {
+        return mm->add_parameter(name, mmlir.get_parameter_shape(name));
+    });
+    std::sort(inputs.begin(), inputs.end(), migraphx::by(std::less<>{}, [](auto ins) {
+                  return to_string(ins->get_operator());
+              }));
+    inputs.push_back(mm->add_parameter("output", mmlir.get_output_shapes().front()));
+
+    migraphx::gpu::context ctx;
+    migraphx::gpu::insert_mlir(*mm, mm->end(), compile_mlir(ctx, mmlir), inputs);
+    return p;
+}
+
+migraphx::parameter_map generate_params(const migraphx::program& p)
+{
+    migraphx::parameter_map m;
+    std::size_t i = 0;
+    for(auto&& x : p.get_parameter_shapes())
+    {
+        // m[x.first] = migraphx::fill_argument(x.second, 1);
+        m[x.first] = migraphx::generate_argument(x.second, i++);
+    }
+    return m;
+}
+
+migraphx::argument run_gpu(migraphx::program p, const migraphx::parameter_map& inputs)
+{
+    mlir_gpu_target t;
+    p.compile(t);
+    migraphx::parameter_map m;
+    for(auto&& input : inputs)
+    {
+        m[input.first] = t.copy_to(input.second);
+    }
+    for(auto&& x : p.get_parameter_shapes())
+    {
+        if(m.count(x.first) == 0)
+        {
+            m[x.first] = t.allocate(x.second);
+        }
+    }
+    return t.copy_from(p.eval(m).front());
+}
+
+migraphx::argument run_ref(migraphx::program p, const migraphx::parameter_map& inputs)
+{
+    p.compile(migraphx::ref::target{});
+    return p.eval(inputs).front();
+}
+
+bool verify_mlir(const migraphx::module& mmlir)
+{
+    migraphx::program ref;
+    ref.get_main_module()->insert_instructions(ref.get_main_module()->end(), &mmlir);
+
+    auto inputs = generate_params(ref);
+
+    auto mlir = create_program_from_mlir(mmlir);
+    return migraphx::verify_args("mlir", run_ref(ref, inputs), run_gpu(mlir, inputs));
+}
+
+TEST_CASE(conv)
+{
+    const std::string mlir_output = R"__migraphx__(
+module {
+  func @main(%arg0: tensor<2x8x3x3xf32>, %arg1: tensor<1x8x4x4xf32>) -> tensor<1x2x2x2xf32> attributes {kernel = "mixr"} {
+    %0 = migraphx.convolution(%arg1, %arg0) {dilation = [1, 1], group = 1 : i64, padding = [0, 0, 0, 0], padding_mode = 0 : i64, stride = [1, 1]} : (tensor<1x8x4x4xf32>, tensor<2x8x3x3xf32>) -> tensor<1x2x2x2xf32>
+    return %0 : tensor<1x2x2x2xf32>
+  }
+}
+)__migraphx__";
+    migraphx::module m;
+    auto x    = m.add_parameter("x", {migraphx::shape::float_type, {1, 8, 4, 4}});
+    auto w    = m.add_parameter("w", {migraphx::shape::float_type, {2, 8, 3, 3}});
+    auto conv = m.add_instruction(migraphx::make_op("convolution"), x, w);
+    m.add_return({conv});
+    auto s = migraphx::gpu::dump_mlir(m);
+    // Skip test if MLIR is not enabled
+    if(s.empty())
+        return;
+    CHECK(encode(s) == encode(mlir_output));
+    EXPECT(verify_mlir(m));
+}
+
+TEST_CASE(conv_add_relu)
+{
+    const std::string mlir_output = R"__migraphx__(
+module {
+  func @main(%arg0: tensor<1x2x2x2xf32>, %arg1: tensor<2x8x3x3xf32>, %arg2: tensor<1x8x4x4xf32>) -> tensor<1x2x2x2xf32> attributes {kernel = "mixr"} {
+    %0 = migraphx.convolution(%arg2, %arg1) {dilation = [1, 1], group = 1 : i64, padding = [0, 0, 0, 0], padding_mode = 0 : i64, stride = [1, 1]} : (tensor<1x8x4x4xf32>, tensor<2x8x3x3xf32>) -> tensor<1x2x2x2xf32>
+    %1 = migraphx.add(%0, %arg0) : (tensor<1x2x2x2xf32>, tensor<1x2x2x2xf32>) -> tensor<1x2x2x2xf32>
+    %2 = migraphx.relu(%1) : (tensor<1x2x2x2xf32>) -> tensor<1x2x2x2xf32>
+    return %2 : tensor<1x2x2x2xf32>
+  }
+}
+)__migraphx__";
+    migraphx::module m;
+    auto x    = m.add_parameter("x", {migraphx::shape::float_type, {1, 8, 4, 4}});
+    auto w    = m.add_parameter("w", {migraphx::shape::float_type, {2, 8, 3, 3}});
+    auto b    = m.add_parameter("b", {migraphx::shape::float_type, {1, 2, 2, 2}});
+    auto conv = m.add_instruction(migraphx::make_op("convolution"), x, w);
+    auto add  = m.add_instruction(migraphx::make_op("add"), conv, b);
+    auto relu = m.add_instruction(migraphx::make_op("relu"), add);
+    m.add_return({relu});
+    auto s = migraphx::gpu::dump_mlir(m);
+    // Skip test if MLIR is not enabled
+    if(s.empty())
+        return;
+    CHECK(encode(s) == encode(mlir_output));
+    EXPECT(verify_mlir(m));
+}
+
+int main(int argc, const char* argv[]) { test::run(argc, argv); }