Enable MLIR by default in migraphx builds (#2037)

.github/workflows/ci.yaml

@@ -274,11 +274,10 @@ jobs:
         # This path is specific to Ubuntu
         path: ${{ github.workspace }}/cget
         # Look to see if there is a cache hit for the corresponding requirements file
-        key: ${{ matrix.os }}-cget-4-${{ hashFiles('requirements.txt', 'dev-requirements.txt') }}
+        key: ${{ matrix.os }}-cget-4-${{ hashFiles('requirements.txt', 'dev-requirements.txt', 'rbuild.ini') }}
         restore-keys: ${{ matrix.os }}-cget-4-
 
     - name: Install dependencies
-      if: steps.deps_cache.outputs.cache-hit != 'true'
       run: rbuild prepare -d cget -s gh
 
     - name: Restore cache files for ccache

Dockerfile

@@ -77,6 +77,9 @@ ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
 ADD rbuild.ini /rbuild.ini
 
+# Temporarily install a new cmake until switching to ubuntu 22.04
+RUN pip3 install cmake==3.22.1
+
 COPY ./tools/install_prereqs.sh /
 RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh
 RUN test -f /usr/local/hash || exit 1
@@ -113,9 +116,6 @@ RUN git clone --single-branch --branch ${ONNXRUNTIME_BRANCH} --recursive ${ONNXR
 
 ADD tools/build_and_test_onnxrt.sh /onnxruntime/build_and_test_onnxrt.sh
 
-# Use the /opt/cmake install because LLVM/MLIR need cmake >= 3.20
-RUN env PATH=/opt/cmake/bin:$PATH cget -p /usr/local install ROCmSoftwarePlatform/rocMLIR@ea15b3597ce55b9088621818228595dd48fb6ec0 -DBUILD_FAT_LIBROCKCOMPILER=On
-
 ENV MIOPEN_FIND_DB_PATH=/tmp/miopen/find-db
 ENV MIOPEN_USER_DB_PATH=/tmp/miopen/user-db
 ENV LD_LIBRARY_PATH=$PREFIX/lib

dev-requirements.txt

@@ -26,5 +26,5 @@ facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cmake
 ccache@v4.1 -DENABLE_TESTING=OFF
 pcre,pfultz2/pcre@8.45 -H sha256:d6f7182602a775a7d500a0cedca6449af0400c6493951513046d17615ed0bf11
 danmar/cppcheck@bb2711c22a0be09efe7f1a8da3030876471026c8 -DHAVE_RULES=1 # 2.11
-RadeonOpenCompute/rocm-cmake@027404a8326da6e7e9338e0b81f9428660190724 --build
+RadeonOpenCompute/rocm-cmake@189d497ed185683154ae9766393b9a10ff21201f --build
 -f requirements.txt

hip-clang.docker

@@ -54,5 +54,8 @@ ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
 ADD rbuild.ini /rbuild.ini
 
+# Temporarily install a new cmake until switching to ubuntu 22.04
+RUN pip3 install cmake==3.22.1
+
 COPY ./tools/install_prereqs.sh /
 RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh

rbuild.ini

@@ -6,7 +6,9 @@ deps =
     -f requirements.txt
 
 [gh]
-ignore = danmar/cppcheck
+ignore =
+    danmar/cppcheck
+    ROCmSoftwarePlatform/rocMLIR
 deps =
     -f dev-requirements.txt
     oneapi-src/oneDNN@v1.7

requirements.txt

@@ -29,3 +29,4 @@ pybind/pybind11@d159a563383d10c821ba7b2a71905d1207db6de4 --build
 msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off
 sqlite3@3.17 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 ROCmSoftwarePlatform/composable_kernel@5172ec5280f14974beee2acf1af1db3b2670244c -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
+ROCmSoftwarePlatform/rocMLIR@ea15b3597ce55b9088621818228595dd48fb6ec0 -DBUILD_FAT_LIBROCKCOMPILER=On

src/targets/gpu/CMakeLists.txt

@@ -191,7 +191,7 @@ register_op(migraphx_gpu HEADER migraphx/gpu/convolution.hpp
 rocm_set_soversion(migraphx_gpu ${MIGRAPHX_SO_VERSION})
 rocm_clang_tidy_check(migraphx_gpu)
 
-set(MIGRAPHX_ENABLE_MLIR OFF CACHE BOOL "")
+set(MIGRAPHX_ENABLE_MLIR ON CACHE BOOL "")
 
 if(MIGRAPHX_ENABLE_MLIR)
     # Find package rocMLIR

src/targets/gpu/fuse_mlir.cpp

@@ -86,7 +86,7 @@ struct mlir_op
         size_t param_cnt               = 0;
         std::vector<std::string> names = mod->get_parameter_names();
         std::sort(names.begin(), names.end());
-        for(std::string param_name : names)
+        for(const std::string& param_name : names)
         {
             ins_shapes[mod->get_parameter(param_name)] = inputs[param_cnt++];
         }
@@ -210,7 +210,8 @@ struct find_mlir_op
             return false;
         }
         const std::initializer_list<std::string> any_type_ops = {"@literal", "@param", "@return"};
-        const std::initializer_list<std::string> no_bool_ops  = {"convolution",
+        const std::initializer_list<std::string> no_bool_ops  = {
+            "convolution",
             "quant_convolution",
             "dot",
             "quant_dot",
@@ -225,17 +226,21 @@ struct find_mlir_op
             "quantizelinear",
             "dequantizelinear",
             "abs",
-                                                                "neg"};
-        const std::initializer_list<std::string> fp_only_ops  = {"ceil",
+            "neg",
+        };
+        const std::initializer_list<std::string> fp_only_ops = {
+            "ceil",
             "erf",
             "exp",
             "floor",
             "log",
             "recip",
             "rsqrt",
-                                                                "sigmoid"
+            // There are bugs in MLIR right now for models using sigmoid so disable it for now
+            // "sigmoid",
             "softmax",
-                                                                "tanh"};
+            "tanh",
+        };
         bool is_float = contains({type_t::float_type, type_t::half_type}, result_type);
         if(contains(any_type_ops, name))
             return true;

src/targets/gpu/mlir.cpp

@@ -644,7 +644,7 @@ struct mlir_program
 
     void set_gpu_properties(const context& migraphx_ctx)
     {
-        auto& device = migraphx_ctx.get_current_device();
+        const auto& device = migraphx_ctx.get_current_device();
         target_arch  = device.get_device_name();
         num_cu       = device.get_cu_count();
     }
@@ -669,10 +669,10 @@ struct mlir_program
         MIGRAPHX_THROW("Failed to compile mlir program");
     }
 
-    void set_tuning(const value& v)
+    void set_tuning(const value& v) MIGRAPHX_TIDY_CONST
     {
-        auto* str = v.if_string();
-        if(not str)
+        const auto* str = v.if_string();
+        if(str == nullptr)
             MIGRAPHX_THROW("mlir tuning solutions must be strings");
         if(not mlirRockTuningSetFromStr(mmodule.get(), make_mlir_string_ref(*str)))
             MIGRAPHX_THROW("Failed setting tuning key: " + *str);
@@ -747,10 +747,10 @@ struct mlir_program
                 {
                     std::vector<std::string> tokens = split_string(line, '\t');
                     std::string arch                = tokens[0];
-                    std::string numCU               = tokens[1];
+                    std::string num_cu              = tokens[1];
                     std::string prob                = tokens[2];
                     std::string perf                = tokens[3];
-                    std::string key = arch.append("\t").append(numCU).append("\t").append(prob);
+                    std::string key = arch.append("\t").append(num_cu).append("\t").append(prob);
                     mlirRockTuningUpdateTable(tuning_table.get(),
                                               make_mlir_string_ref(key),
                                               make_mlir_string_ref(perf),