Merge branch 'develop' into lw/q_linear_add

.github/workflows/ci.yaml

@@ -137,7 +137,6 @@ jobs:
           -DMIGRAPHX_ENABLE_GPU=On \
           -DMIGRAPHX_ENABLE_CPU=On \
           -DMIGRAPHX_ENABLE_FPGA=On \
-          -DMIGRAPHX_ENABLE_MLIR=On \
           -DBUILD_DEV=On \
           -DROCM_ENABLE_GH_ANNOTATIONS=On \
           -DCLANG_TIDY_DEPEND_ON_TARGET=Off \
@@ -282,6 +281,7 @@ jobs:
           -DBUILD_DEV=On \
           -DCMAKE_CXX_COMPILER_LAUNCHER=/usr/local/bin/ccache \
           -DCMAKE_C_COMPILER_LAUNCHER=/usr/local/bin/ccache \
+          -DCMAKE_CXX_FLAGS="-Werror" \
           -DGPU_TARGETS=gfx908 \
           ..
         make -j$(nproc) tests driver

CMakeLists.txt

@@ -71,7 +71,7 @@ include(ROCMSetupVersion)
 
 option(BUILD_DEV "Build for development purpose only" OFF)
 
-rocm_setup_version(VERSION 2.7.0)
+rocm_setup_version(VERSION 2.8.0)
 set(MIGRAPHX_SO_VERSION ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH})
 
 option( BUILD_SHARED_LIBS "Build as a shared library" ON )
@@ -283,7 +283,6 @@ file(MAKE_DIRECTORY ${DEST_DIR}/lib/onnx_migraphx)
 foreach(py_file ${backend_files})
     configure_file(${py_file} ${DEST_DIR}/lib/onnx_migraphx/. COPYONLY)
 endforeach(py_file)
-configure_file(${CMAKE_SOURCE_DIR}/test/py/onnx_backend_test.py ${DEST_DIR}/onnx_backend_test.py COPYONLY)
 
 rocm_create_package(
     NAME MIGraphX

Dockerfile

@@ -101,10 +101,6 @@ RUN cget -p $PREFIX install facebook/[email protected] -X subdir -DCMAKE_DIR=build/cma
 RUN cget -p $PREFIX install [email protected] -DENABLE_TESTING=OFF
 RUN cget -p /opt/cmake install kitware/[email protected]
 
-# Install MLIR
-ADD mlir-requirements.txt /mlir-requirements.txt
-RUN cget -p /usr/local install -f /mlir-requirements.txt
-
 COPY ./test/onnx/.onnxrt-commit /
 
 ARG ONNXRUNTIME_REPO=https://github.com/Microsoft/onnxruntime

Jenkinsfile

@@ -1,3 +1,7 @@
+def getgputargets() {
+    targets="gfx906;gfx908;gfx90a;gfx1030;gfx1100;gfx1101;gfx1102"
+    return targets
+}
 
 // def rocmtestnode(variant, name, body, args, pre) {
 def rocmtestnode(Map conf) {
@@ -107,11 +111,13 @@ rocmtest clang_debug: rocmnode('cdna') { cmake_build ->
     stage('hipRTC Debug') {
         def sanitizers = "undefined"
         def debug_flags = "-g -O2 -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
-        cmake_build(flags: "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}' -DCMAKE_C_FLAGS_DEBUG='${debug_flags}' -DMIGRAPHX_USE_HIPRTC=On -DGPU_TARGETS=$(/opt/rocm/bin/rocminfo | grep -o -m1 'gfx.*')", gpu_debug: true)
+        def gpu_targets = getgputargets()
+        cmake_build(flags: "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}' -DCMAKE_C_FLAGS_DEBUG='${debug_flags}' -DMIGRAPHX_USE_HIPRTC=On -DGPU_TARGETS='${gpu_targets}'", gpu_debug: true)
     }
 }, clang_release: rocmnode('mi100+') { cmake_build ->
     stage('Hip Clang Release') {
-        cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DGPU_TARGETS=$(/opt/rocm/bin/rocminfo | grep -o -m1 'gfx.*')")
+        def gpu_targets = getgputargets()
+        cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DGPU_TARGETS='${gpu_targets}'")
         stash includes: 'build/*.deb', name: 'migraphx-package'
     }
 // }, hidden_symbols: rocmnode('cdna') { cmake_build ->
@@ -120,7 +126,8 @@ rocmtest clang_debug: rocmnode('cdna') { cmake_build ->
 //     }
 }, all_targets_debug : rocmnode('cdna') { cmake_build ->
     stage('All targets Release') {
-        cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DMIGRAPHX_ENABLE_GPU=On -DMIGRAPHX_ENABLE_CPU=On -DMIGRAPHX_ENABLE_FPGA=On -DGPU_TARGETS=$(/opt/rocm/bin/rocminfo | grep -o -m1 'gfx.*')") 
+        def gpu_targets = getgputargets()
+        cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DMIGRAPHX_ENABLE_GPU=On -DMIGRAPHX_ENABLE_CPU=On -DMIGRAPHX_ENABLE_FPGA=On -DGPU_TARGETS='${gpu_targets}'")
     }
 }, mlir_debug: rocmnode('cdna') { cmake_build ->
     stage('MLIR Debug') {
@@ -129,20 +136,23 @@ rocmtest clang_debug: rocmnode('cdna') { cmake_build ->
             // Note: the -fno-sanitize= is copied from upstream LLVM_UBSAN_FLAGS.
             def debug_flags_cxx = "-g -O2 -fsanitize=${sanitizers} -fno-sanitize=vptr,function -fno-sanitize-recover=${sanitizers}"
             def debug_flags = "-g -O2 -fsanitize=${sanitizers} -fno-sanitize=vptr -fno-sanitize-recover=${sanitizers}"
-            cmake_build(flags: "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DMIGRAPHX_ENABLE_MLIR=On -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags_cxx}' -DCMAKE_C_FLAGS_DEBUG='${debug_flags}' -DGPU_TARGETS=$(/opt/rocm/bin/rocminfo | grep -o -m1 'gfx.*')")
+            def gpu_targets = getgputargets()
+            cmake_build(flags: "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DMIGRAPHX_ENABLE_MLIR=On -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags_cxx}' -DCMAKE_C_FLAGS_DEBUG='${debug_flags}' -DGPU_TARGETS='${gpu_targets}'")
         }
     }
-}, ck_release: rocmnode('mi100+') { cmake_build ->
-    stage('CK Release') {
+}, ck_hiprtc: rocmnode('mi100+') { cmake_build ->
+    stage('CK hipRTC') {
         withEnv(['MIGRAPHX_ENABLE_CK=1', 'MIGRAPHX_TUNE_CK=1']) {
-            cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DGPU_TARGETS=$(/opt/rocm/bin/rocminfo | grep -o -m1 'gfx.*')")
+            def gpu_targets = getgputargets()
+            cmake_build(flags: "-DCMAKE_BUILD_TYPE=release -DMIGRAPHX_USE_HIPRTC=On -DGPU_TARGETS='${gpu_targets}'")
         }
     }
 }, clang_asan: rocmnode('nogpu') { cmake_build ->
     stage('Clang ASAN') {
         def sanitizers = "undefined,address"
         def debug_flags = "-g -O2 -fno-omit-frame-pointer -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
-        cmake_build(flags: "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DMIGRAPHX_ENABLE_GPU=Off -DMIGRAPHX_ENABLE_CPU=On -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}' -DCMAKE_C_FLAGS_DEBUG='${debug_flags}'")
+        def gpu_targets = getgputargets()
+        cmake_build(flags: "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DMIGRAPHX_ENABLE_GPU=Off -DMIGRAPHX_ENABLE_CPU=On -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}' -DCMAKE_C_FLAGS_DEBUG='${debug_flags}' -DGPU_TARGETS='${gpu_targets}'")
     }
 }//, clang_release_navi: rocmnode('navi21') { cmake_build ->
 //    stage('HIP Clang Release Navi') {

cmake/Embed.cmake

@@ -77,16 +77,17 @@ function(generate_embed_source EMBED_NAME)
         list(GET PARSE_FILES ${idx} FILE)
 
         set(START_SYMBOL "_binary_${SYMBOL}_start")
-        set(END_SYMBOL "_binary_${SYMBOL}_end")
+        set(LENGTH_SYMBOL "_binary_${SYMBOL}_length")
         if(EMBED_USE_LD)
             string(APPEND EXTERNS "
-                extern const char ${START_SYMBOL}[];
-                extern const char ${END_SYMBOL}[];
+extern const char ${START_SYMBOL}[];
+extern const size_t _binary_${SYMBOL}_size;
+const auto ${LENGTH_SYMBOL} = reinterpret_cast<size_t>(&_binary_${SYMBOL}_size);
             ")
         else()
             string(APPEND EXTERNS "
-                extern const char ${START_SYMBOL}[];
-                extern const char* ${END_SYMBOL};
+extern const char ${START_SYMBOL}[];
+extern const size_t ${LENGTH_SYMBOL};
             ")
         endif()
 
@@ -97,23 +98,22 @@ function(generate_embed_source EMBED_NAME)
         endif()
 
         string(APPEND INIT_KERNELS "
-            { \"${BASE_NAME}\", { ${START_SYMBOL}, ${END_SYMBOL}} },
-        ")
+            { \"${BASE_NAME}\", { ${START_SYMBOL}, ${LENGTH_SYMBOL}} },")
     endforeach()
 
     file(WRITE "${PARSE_HEADER}" "
+#include <string_view>
 #include <unordered_map>
-#include <string>
 #include <utility>
-const std::unordered_map<std::string, std::pair<const char*,const char*>>& ${EMBED_NAME}();
+std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}();
 ")
 
     file(WRITE "${PARSE_SRC}" "
 #include <${EMBED_NAME}.hpp>
 ${EXTERNS}
-const std::unordered_map<std::string, std::pair<const char*,const char*>>& ${EMBED_NAME}()
+std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}()
 {
-    static const std::unordered_map<std::string, std::pair<const char*,const char*>> result = {${INIT_KERNELS}};
+    static std::unordered_map<std::string_view, std::string_view> result = {${INIT_KERNELS}};
     return result;
 }
 ")
@@ -154,9 +154,10 @@ function(embed_file OUTPUT_FILE OUTPUT_SYMBOL FILE)
             # removes trailing comma
             string(REGEX REPLACE ", $" "" ARRAY_VALUES ${ARRAY_VALUES})
             file(WRITE "${OUT_FILE}" "
-                extern const char _binary_${SYMBOL}_start[] = { ${ARRAY_VALUES} };
-                extern const char* _binary_${SYMBOL}_end = _binary_${SYMBOL}_start + sizeof(_binary_${SYMBOL}_start);
-            \n")
+#include <cstddef>
+extern const char _binary_${SYMBOL}_start[] = { ${ARRAY_VALUES} };
+extern const size_t _binary_${SYMBOL}_length = sizeof(_binary_${SYMBOL}_start);
+")
         endif()
     endforeach()
 endfunction()

hip-clang.docker

@@ -60,6 +60,3 @@ RUN pip3 install cmake==3.22.1
 COPY ./tools/install_prereqs.sh /
 RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh
 
-# Install MLIR
-ADD mlir-requirements.txt /mlir-requirements.txt
-RUN cget -p /usr/local install -f /mlir-requirements.txt

requirements.txt

@@ -28,4 +28,5 @@ ROCmSoftwarePlatform/[email protected]
 pybind/pybind11@d159a563383d10c821ba7b2a71905d1207db6de4 --build
 msgpack/[email protected] -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/composable_kernel@a22e479b8e1557961039db2d5c5ff89cff35e86b -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
+ROCmSoftwarePlatform/rocMLIR@12748a3402c069f733ea7f2ba1f8d8a070b3622a -DBUILD_FAT_LIBROCKCOMPILER=On

src/compile_src.cpp

@@ -46,7 +46,7 @@ std::vector<char> src_compiler::compile(const std::vector<src_file>& srcs) const
         fs::path full_path   = td.path / src.path;
         fs::path parent_path = full_path.parent_path();
         fs::create_directories(parent_path);
-        write_buffer(full_path.string(), src.content.first, src.len());
+        write_buffer(full_path.string(), src.content.data(), src.content.size());
         if(src.path.extension().string() == ".cpp")
         {
             params += " " + src.path.filename().string();

src/cpp_generator.cpp

@@ -213,13 +213,13 @@ cpp_generator::function cpp_generator::generate_module(const module& m,
                 ins->get_literal().visit([&](auto v) {
                     assert(v.size() == 1);
                     auto x = v.front();
-                    if(std::isinf(x))
+                    if(std::isinf(static_cast<double>(x)))
                     {
                         string_literal = "__builtin_huge_val()";
                         if(x < 0)
                             string_literal = "-__builtin_huge_val()";
                     }
-                    else if(std::isnan(x))
+                    else if(std::isnan(static_cast<double>(x)))
                         string_literal = "__builtin_nan()";
                     else
                         string_literal = ins->get_literal().to_string();

src/include/migraphx/compile_src.hpp

@@ -37,8 +37,18 @@ inline namespace MIGRAPHX_INLINE_NS {
 struct src_file
 {
     fs::path path;
-    std::pair<const char*, const char*> content;
-    std::size_t len() const { return content.second - content.first; }
+    std::string_view content;
+
+    src_file() = default;
+    src_file(fs::path file_path, std::string_view file_content)
+        : path{std::move(file_path)}, content{file_content}
+    {
+    }
+
+    explicit src_file(const std::pair<std::string_view, std::string_view>& pair)
+        : path{pair.first}, content{pair.second}
+    {
+    }
 };
 
 struct MIGRAPHX_EXPORT src_compiler

src/include/migraphx/normalize_attributes.hpp

@@ -52,6 +52,7 @@ using dependent_type = typename select_dependent_type<T, Ts...>::type;
  * \param attr_val the normalize_axes attributes from the operator
  * \param prefix error message prefix
  */
+MIGRAPHX_EXPORT
 std::vector<int64_t> normalize_axes(const std::vector<int64_t>& axes,
                                     const shape& input_shape,
                                     const value& attr_val,
@@ -67,6 +68,7 @@ std::vector<int64_t> normalize_axes(const std::vector<int64_t>& axes,
  * \param attr_val the normalize_axes attributes from the operator
  * \param prefix error message prefix
  */
+MIGRAPHX_EXPORT
 std::vector<int64_t> normalize_indices(const std::vector<int64_t>& indices,
                                        const std::vector<int64_t>& axes,
                                        const shape& input_shape,

src/include/migraphx/op/allocate.hpp

@@ -33,6 +33,19 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
+/**
+ * Static allocate:
+ * No inputs: `allocate()`
+ * `this.s` attribute set to the static output shape of the buffer.
+ *
+ * Dynamic allocate:
+ * One input: `allocate(output_dims)`
+ * `output_dims` are the output buffer dimensions and has a static shape.
+ * Either `this.s` or `this.buf_type` must be set to calculate the dynamic output shape at compute
+ * time. If `this.buf_type` is set, the compute_shape() of allocate at compile time will have
+ * dynamic_dimensions from {0, max_int} with rank = output_dims.ndim(). If `this.s` is set then the
+ * compute_shape() will output `this.s`; `this.s` should be a dynamic shape.
+ */
 struct allocate
 {
     shape s{};

src/include/migraphx/op/convert.hpp

@@ -68,7 +68,7 @@ struct convert : unary<convert>
             auto y = x;
             shape::visit(type, [&](auto as) {
                 // clamping value between target_type's max and min doesn't work for NaNs,
-                if(std::isnan(x))
+                if(std::isnan(static_cast<double>(x)))
                 {
                     y = as.nan();
                 }

src/include/migraphx/op/isnan.hpp

@@ -35,7 +35,7 @@ struct isnan : unary<isnan>
 {
     auto apply() const
     {
-        return [](auto x) { return std::isnan(x); };
+        return [](auto x) { return std::isnan(static_cast<double>(x)); };
     }
 
     std::string name() const { return "isnan"; }

src/include/migraphx/pad_calc.hpp

@@ -64,6 +64,7 @@ shape compute_padded_shape(const shape& input,
 
 // Used for dynamic auto padding of pooling operators where padding needs to be computed at
 // evaulation time.
+MIGRAPHX_EXPORT
 shape compute_padded_pool_shape(const shape& input,
                                 const shape& kernel,
                                 const std::vector<std::size_t>& padding,

src/include/migraphx/verify.hpp

@@ -90,8 +90,7 @@ struct not_finite_fn
     template <class T>
     bool operator()(T x) const
     {
-        using std::isfinite;
-        return not isfinite(x);
+        return not std::isfinite(static_cast<double>(x));
     }
 };
 static constexpr not_finite_fn not_finite{};
@@ -101,8 +100,7 @@ struct compare_mag_fn
     template <class T, class U>
     bool operator()(T x, U y) const
     {
-        using std::fabs;
-        return fabs(x) < fabs(y);
+        return std::fabs(x) < std::fabs(y);
     }
 };
 static constexpr compare_mag_fn compare_mag{};

src/onnx/parse_castlike.cpp

@@ -0,0 +1,53 @@
+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2023 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/onnx/op_parser.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/instruction.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace onnx {
+
+struct parse_castlike : op_parser<parse_castlike>
+{
+    std::vector<op_desc> operators() const { return {{"CastLike"}}; }
+
+    instruction_ref parse(const op_desc& /*opd*/,
+                          const onnx_parser& /*parser*/,
+                          const onnx_parser::node_info& info,
+                          const std::vector<instruction_ref>& args) const
+    {
+        if(not(args.size() == 2))
+        {
+            MIGRAPHX_THROW("PARSE_CASTLIKE: CastLike must have exactly 2 inputs!");
+        }
+        shape::type_t target_type = args[1]->get_shape().type();
+        return info.add_instruction(make_op("convert", {{"target_type", target_type}}), args[0]);
+    }
+};
+
+} // namespace onnx
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx