Merge branch 'develop' into enable_navi_32_ci

test/include/test.hpp

@@ -22,6 +22,7 @@
  * THE SOFTWARE.
  */
 
+#include <atomic>
 #include <algorithm>
 #include <cassert>
 #include <cstdio>
@@ -342,11 +343,19 @@ inline std::ostream& operator<<(std::ostream& os, const color& c)
     return os;
 }
 
+inline std::atomic<int>& failures()
+{
+    // NOLINTNEXTLINE
+    static std::atomic<int> f = 0;
+    return f;
+}
+
 template <class T, class F>
 void failed(T x, const char* msg, const char* func, const char* file, int line, F f)
 {
     if(not bool(x.value()))
     {
+        failures()++;
         std::cout << func << std::endl;
         std::cout << file << ":" << line << ":" << std::endl;
         std::cout << color::bold << color::fg_red << "    FAILED: " << color::reset << msg << " "
@@ -586,13 +595,21 @@ struct driver
         {
             try
             {
+                failures() = 0;
                 f();
             }
+            // cppcheck-suppress EmptyCatchStatement
             catch(const failure_error&)
             {
-                msg = "Test failure";
             }
         }
+        if(msg.empty() and failures() != 0)
+        {
+            if(failures() == 1)
+                msg = "Test failure";
+            else
+                msg = std::to_string(failures()) + " test failures";
+        }
         if(msg.empty())
         {
             out() << color::fg_green << "[ COMPLETE ] " << color::reset << color::bold << name
@@ -683,10 +700,10 @@ inline void run(int argc, const char* argv[])
 #define TEST_CAPTURE(...) test::capture{}->*__VA_ARGS__
 
 // NOLINTNEXTLINE
-#define CHECK(...)                                                                                 \
-    test::failed(                                                                                  \
-        test::capture{}->*__VA_ARGS__, #__VA_ARGS__, __PRETTY_FUNCTION__, __FILE__, __LINE__, [] { \
-        })
+#define CHECK(...) \
+    test::failed(  \
+        TEST_CAPTURE(__VA_ARGS__), #__VA_ARGS__, __PRETTY_FUNCTION__, __FILE__, __LINE__, [] {})
+
 // NOLINTNEXTLINE
 #define EXPECT(...)                         \
     test::failed(TEST_CAPTURE(__VA_ARGS__), \

test/memory_coloring_test.cpp

@@ -89,17 +89,13 @@ bool is_overlap_load(migraphx::instruction_ref a, migraphx::instruction_ref b)
 
 bool is_disjoint(const std::vector<migraphx::instruction_ref>& inss)
 {
-    for(auto ins1 : inss)
-    {
-        for(auto ins2 : inss)
-        {
+    return std::none_of(inss.begin(), inss.end(), [&](auto ins1) {
+        return std::none_of(inss.begin(), inss.end(), [&](auto ins2) {
             if(ins1 == ins2)
-                continue;
-            if(is_overlap_load(ins1, ins2))
-                return false;
-        }
-    }
-    return true;
+                return true;
+            return is_overlap_load(ins1, ins2);
+        });
+    });
 }
 
 TEST_CASE(test1)

test/onnx/.onnxrt-commit

-21a71d52bd2074b770807b209939ec11e2c64fa7
+e5bb7aba502f5a8783de945258d226c092c14386

test/onnx/onnx_test.cpp

@@ -4712,14 +4712,16 @@ TEST_CASE(quantizelinear_test)
     auto l1  = mm->add_parameter("1", {migraphx::shape::float_type, {1}});
     auto l1_mbcast =
         mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {5}}}), l1);
-    auto div   = mm->add_instruction(migraphx::make_op("div"), l0, l1_mbcast);
-    auto round = mm->add_instruction(migraphx::make_op("round"), div);
-    auto s     = round->get_shape();
-    std::vector<int> min_data(s.elements(), 0);
-    std::vector<int> max_data(s.elements(), 255);
-    auto min_arg = mm->add_literal(s, min_data);
-    auto max_arg = mm->add_literal(s, max_data);
-    auto clip    = mm->add_instruction(migraphx::make_op("clip"), round, min_arg, max_arg);
+    auto div     = mm->add_instruction(migraphx::make_op("div"), l0, l1_mbcast);
+    auto round   = mm->add_instruction(migraphx::make_op("round"), div);
+    auto s       = round->get_shape();
+    auto min_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {0}});
+    auto max_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {255}});
+    auto min_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), min_arg);
+    auto max_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), max_arg);
+    auto clip = mm->add_instruction(migraphx::make_op("clip"), round, min_mbcast, max_mbcast);
     mm->add_instruction(
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::uint8_type)}}),
@@ -4741,14 +4743,16 @@ TEST_CASE(quantizelinear_int32_test)
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::float_type)}}),
         l0);
-    auto div   = mm->add_instruction(migraphx::make_op("div"), l0, l1_mbcast);
-    auto round = mm->add_instruction(migraphx::make_op("round"), div);
-    auto s     = round->get_shape();
-    std::vector<int> min_data(s.elements(), 0);
-    std::vector<int> max_data(s.elements(), 255);
-    auto min_arg = mm->add_literal(s, min_data);
-    auto max_arg = mm->add_literal(s, max_data);
-    auto clip    = mm->add_instruction(migraphx::make_op("clip"), round, min_arg, max_arg);
+    auto div     = mm->add_instruction(migraphx::make_op("div"), l0, l1_mbcast);
+    auto round   = mm->add_instruction(migraphx::make_op("round"), div);
+    auto s       = round->get_shape();
+    auto min_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {0}});
+    auto max_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {255}});
+    auto min_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), min_arg);
+    auto max_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), max_arg);
+    auto clip = mm->add_instruction(migraphx::make_op("clip"), round, min_mbcast, max_mbcast);
     mm->add_instruction(
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::uint8_type)}}),
@@ -4775,13 +4779,15 @@ TEST_CASE(quantizelinear_zero_point_test)
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::float_type)}}),
         l2_mbcast);
-    auto add = mm->add_instruction(migraphx::make_op("add"), round, l2_mbcast);
-    auto s   = round->get_shape();
-    std::vector<int> min_data(s.elements(), -128);
-    std::vector<int> max_data(s.elements(), 127);
-    auto min_arg = mm->add_literal(s, min_data);
-    auto max_arg = mm->add_literal(s, max_data);
-    auto clip    = mm->add_instruction(migraphx::make_op("clip"), add, min_arg, max_arg);
+    auto add     = mm->add_instruction(migraphx::make_op("add"), round, l2_mbcast);
+    auto s       = round->get_shape();
+    auto min_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {-128}});
+    auto max_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {127}});
+    auto min_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), min_arg);
+    auto max_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), max_arg);
+    auto clip = mm->add_instruction(migraphx::make_op("clip"), add, min_mbcast, max_mbcast);
     mm->add_instruction(
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::int8_type)}}),
@@ -4812,13 +4818,15 @@ migraphx::program make_quantizelinear_axis_prog()
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::float_type)}}),
         l2_bcast);
-    auto add = mm->add_instruction(migraphx::make_op("add"), round, l2_bcast);
-    auto s   = round->get_shape();
-    std::vector<int> min_data(s.elements(), -128);
-    std::vector<int> max_data(s.elements(), 127);
-    auto min_arg = mm->add_literal(s, min_data);
-    auto max_arg = mm->add_literal(s, max_data);
-    auto clip    = mm->add_instruction(migraphx::make_op("clip"), add, min_arg, max_arg);
+    auto add     = mm->add_instruction(migraphx::make_op("add"), round, l2_bcast);
+    auto s       = round->get_shape();
+    auto min_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {-128}});
+    auto max_arg = mm->add_literal(migraphx::literal{migraphx::shape{s.type()}, {127}});
+    auto min_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), min_arg);
+    auto max_mbcast =
+        mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), max_arg);
+    auto clip = mm->add_instruction(migraphx::make_op("clip"), add, min_mbcast, max_mbcast);
     mm->add_instruction(
         migraphx::make_op("convert",
                           {{"target_type", migraphx::to_value(migraphx::shape::int8_type)}}),

test/rewrite_quantization_test.cpp

@@ -37,6 +37,17 @@
 
 bool is_quantizelinear(migraphx::instruction& ins) { return ins.name() == "quantizelinear"; }
 bool is_dequantizelinear(migraphx::instruction& ins) { return ins.name() == "dequantizelinear"; }
+bool is_clip_scalar(migraphx::instruction& ins)
+{
+    if(ins.name() == "clip")
+    {
+        assert(ins.inputs().size() > 1);
+        return (std::all_of(ins.inputs().begin() + 1, ins.inputs().end(), [](auto input) {
+            return input->get_shape().scalar();
+        }));
+    }
+    return false;
+}
 
 void run_pass(migraphx::module& m) { migraphx::run_passes(m, {migraphx::rewrite_quantization{}}); }
 
@@ -70,6 +81,8 @@ TEST_CASE(quantizelinear)
     EXPECT(eval(p1) == eval(p2));
     EXPECT(any_of(*p1.get_main_module(), &is_quantizelinear));
     EXPECT(none_of(*p2.get_main_module(), &is_quantizelinear));
+    // ensure clip literals created in quantized program are scalar
+    EXPECT(any_of(*p2.get_main_module(), &is_clip_scalar));
 }
 
 TEST_CASE(dequantizelinear)

test/simplify_algebra_test.cpp

@@ -2189,16 +2189,16 @@ TEST_CASE(simplify_split_between_add)
     EXPECT(m1.sort() == m2.sort());
 }
 
-TEST_CASE(simplify_dot_horiz)
+void test_dot_horiz(migraphx::shape::type_t type, const std::string& dot_type)
 {
-    auto s = migraphx::shape{migraphx::shape::int32_type, {3, 2, 2}};
+    auto s = migraphx::shape{type, {3, 2, 2}};
     migraphx::module m1;
     {
         auto input = m1.add_parameter("input", s);
         auto a     = m1.add_literal(migraphx::generate_literal(s, 0));
         auto b     = m1.add_literal(migraphx::generate_literal(s, 1));
-        auto x     = m1.add_instruction(migraphx::make_op("dot"), input, a);
-        auto y     = m1.add_instruction(migraphx::make_op("dot"), input, b);
+        auto x     = m1.add_instruction(migraphx::make_op(dot_type), input, a);
+        auto y     = m1.add_instruction(migraphx::make_op(dot_type), input, b);
         auto sum   = m1.add_instruction(migraphx::make_op("add"), x, y);
         m1.add_instruction(pass_op{}, sum);
     }
@@ -2210,7 +2210,7 @@ TEST_CASE(simplify_dot_horiz)
         auto a      = m2.add_literal(migraphx::generate_literal(s, 0));
         auto b      = m2.add_literal(migraphx::generate_literal(s, 1));
         auto concat = m2.add_instruction(migraphx::make_op("concat", {{"axis", 2}}), a, b);
-        auto dot    = m2.add_instruction(migraphx::make_op("dot"), input, concat);
+        auto dot    = m2.add_instruction(migraphx::make_op(dot_type), input, concat);
         auto x      = m2.add_instruction(
             migraphx::make_op("slice", {{"axes", {2}}, {"starts", {0}}, {"ends", {2}}}), dot);
         auto y = m2.add_instruction(
@@ -2221,6 +2221,10 @@ TEST_CASE(simplify_dot_horiz)
     EXPECT(m1.sort() == m2.sort());
 }
 
+TEST_CASE(simplify_dot_horiz) { test_dot_horiz(migraphx::shape::int32_type, "dot"); }
+
+TEST_CASE(simplify_quant_dot_horiz) { test_dot_horiz(migraphx::shape::int8_type, "quant_dot"); }
+
 TEST_CASE(simplify_dot_horiz_same_constant)
 {
     auto s = migraphx::shape{migraphx::shape::int32_type, {3, 2, 2}};

tools/docker/sles.docker

+FROM registry.suse.com/suse/sle15:15.4
+
+RUN sh -c 'echo -e "\
+[rocm]\n\
+name=rocm\n\
+baseurl=https://repo.radeon.com/rocm/zyp/5.5/main\n\
+enabled=1\n\
+gpgcheck=1\n\
+gpgkey=https://repo.radeon.com/rocm/rocm.gpg.key\n\
+" > /etc/zypp/repos.d/rocm.repo'
+
+RUN cat /etc/zypp/repos.d/rocm.repo
+
+RUN zypper -n --gpg-auto-import-keys refresh
+
+RUN zypper install -y -t pattern devel_basis enhanced_base
+RUN zypper --gpg-auto-import-keys install -y \
+    doxygen \
+    gcc-c++ \
+    gdb \
+    git \
+    python3-pip
+
+# Workaround broken rocm packages
+RUN ln -s /opt/rocm-* /opt/rocm
+RUN echo "/opt/rocm/lib" > /etc/ld.so.conf.d/rocm.conf
+RUN echo "/opt/rocm/llvm/lib" > /etc/ld.so.conf.d/rocm-llvm.conf
+RUN ldconfig
+
+ENV LC_ALL=C.UTF-8
+ENV LANG=C.UTF-8
+
+# Install yapf
+RUN pip3 install yapf==0.28.0
+
+# Install doc requirements
+# ADD docs/.sphinx/requirements.txt /doc-requirements.txt
+# RUN pip3 install -r /doc-requirements.txt
+
+# Install dependencies
+ADD dev-requirements.txt /dev-requirements.txt
+ADD requirements.txt /requirements.txt
+ADD rbuild.ini /rbuild.ini
+
+COPY ./tools/install_prereqs.sh /
+RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh

tools/install_prereqs.sh

@@ -31,9 +31,30 @@ set -e
 export LC_ALL=C.UTF-8
 export LANG=C.UTF-8
 
+source /etc/os-release
+
+if [[ ("${ID}" == "sles") ]]; then
+  zypper -n --gpg-auto-import-keys install -y \
+    cmake \
+    miopen-hip-devel \
+    openmp-extras-devel \
+    python3-devel \
+    python3-pip \
+    rocblas-devel \
+    rocm-cmake
+else
+  # Need pip3 and Python headers to build dependencies
+  apt update && apt install -y \
+    cmake \
+    libnuma-dev \
+    miopen-hip-dev \
+    openmp-extras \
+    python3-dev \
+    python3-pip \
+    rocblas-dev \
+    rocm-cmake
+fi
 
-# Need pip3 and Python headers to build dependencies
-apt update && apt install -y python3-pip python3-dev cmake rocm-cmake rocblas miopen-hip openmp-extras
 
 # Needed for cmake to build various pip packages
 pip3 install setuptools wheel
@@ -56,9 +77,11 @@ echo "Dependencies are installed at $PREFIX"
 # Install deps with rbuild
 rbuild prepare -d $PREFIX -s develop
 
+if [[ ("${ID}" != "sles") ]]; then
 export CMAKE_ARGS="-DONNX_USE_PROTOBUF_SHARED_LIBS=ON"
 
 pip3 install onnx==1.10.2 numpy==1.21.6 typing==3.7.4 pytest==6.0.1 packaging==23.0
 
 # pin version of protobuf in Python for onnx runtime unit tests between dist versions
 pip3 install protobuf==3.20.0
+fi