Upgrade docker to rocm 4.1 and drop hcc (#795)

* Fix tidy warnings for 4.1

* Formatting

* Upgrade to 4.1 in docker

* Remove hcc build and enable ubsan on clang debug

* Add missing openmp package

* Construct directly

* Construct directly

* Upgrade rocm-cmake version

CMakeLists.txt

@@ -92,6 +92,7 @@ rocm_enable_clang_tidy(
         # Disable the aliased reserved identifiers
         -cert-dcl37-c
         -cert-dcl51-cpp
+        -cert-str34-c
         # Disable all alpha checks by default
         -clang-analyzer-alpha*
         # Enable some alpha checks
@@ -136,6 +137,7 @@ rocm_enable_clang_tidy(
         -readability-braces-around-statements
         -readability-convert-member-functions-to-static
         -readability-else-after-return
+        -readability-function-cognitive-complexity
         -readability-named-parameter
         -readability-redundant-string-init
         -readability-uppercase-literal-suffix

Dockerfile

@@ -6,7 +6,7 @@ ARG PREFIX=/usr/local
 RUN dpkg --add-architecture i386
 
 # Add rocm repository
-RUN sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/.apt_3.7/ xenial main > /etc/apt/sources.list.d/rocm.list'
+RUN sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/4.1/ xenial main > /etc/apt/sources.list.d/rocm.list'
 
 # Install dependencies
 RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \

Jenkinsfile

@@ -83,17 +83,10 @@ def rocmnode(name, body) {
     }
 }
 
-def rochccmnode(name, body) {
-    return { label ->
-        rocmtestnode(variant: label, node: rocmnodename(name), docker_build_args: '-f hcc.docker', body: body)
-    }
-}
-
 rocmtest clang_debug: rocmnode('vega') { cmake_build ->
     stage('Hip Clang Debug') {
-        // def sanitizers = "undefined"
-        // def debug_flags = "-O2 -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
-        def debug_flags = "-g -O2"
+        def sanitizers = "undefined"
+        def debug_flags = "-g -O2 -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
         cmake_build("/opt/rocm/llvm/bin/clang++", "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}'")
     }
 }, clang_release: rocmnode('vega') { cmake_build ->
@@ -101,13 +94,6 @@ rocmtest clang_debug: rocmnode('vega') { cmake_build ->
         cmake_build("/opt/rocm/llvm/bin/clang++", "-DCMAKE_BUILD_TYPE=release")
         stash includes: 'build/*.deb', name: 'migraphx-package'
     }
-}, hcc_debug: rochccmnode('vega') { cmake_build ->
-    stage('Hcc Debug') {
-        // TODO: Enable integer
-        def sanitizers = "undefined"
-        def debug_flags = "-O2 -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
-        cmake_build("/opt/rocm/bin/hcc", "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}'")
-    }
 }
 
 def onnxnode(name, body) {

hcc.docker

-FROM ubuntu:18.04
-
-ARG PREFIX=/usr/local
-
-# Support multiarch
-RUN dpkg --add-architecture i386
-
-# Add rocm repository
-RUN sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/3.0/ xenial main > /etc/apt/sources.list.d/rocm.list'
-
-# Install dependencies
-RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \
-    apt-utils \
-    build-essential \
-    clang-format-5.0 \
-    cmake \
-    curl \
-    doxygen \
-    gdb \
-    git \
-    lcov \
-    pkg-config \
-    python \
-    python-dev \
-    python-pip \
-    python3 \
-    python3-dev \
-    python3-pip \
-    software-properties-common \
-    wget \
-    rocm-clang-ocl \
-    rocm-device-libs \
-    rocblas \
-    zlib1g-dev && \
-    apt-get clean && \
-    rm -rf /var/lib/apt/lists/*
-
-# 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 rbuild
-RUN pip3 install https://github.com/RadeonOpenCompute/rbuild/archive/master.tar.gz
-
-# Install doc requirements
-ADD doc/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
-# Manually ignore rocm dependencies
-RUN cget -p $PREFIX ignore \
-    RadeonOpenCompute/clang-ocl \
-    ROCm-Developer-Tools/HIP \
-    ROCmSoftwarePlatform/MIOpenGEMM \
-    ROCmSoftwarePlatform/rocBLAS
-RUN cget -p $PREFIX init --cxx /opt/rocm/bin/hcc
-RUN cget -p $PREFIX install -f dev-requirements.txt
-
-# Install latest ccache version
-RUN cget -p $PREFIX install facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cmake
-RUN cget -p $PREFIX install ccache@v4.1
-

hip-clang.docker

@@ -6,7 +6,7 @@ ARG PREFIX=/usr/local
 RUN dpkg --add-architecture i386
 
 # Add rocm repository
-RUN sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/.apt_3.7/ xenial main > /etc/apt/sources.list.d/rocm.list'
+RUN sh -c 'echo deb [arch=amd64 trusted=yes] http://repo.radeon.com/rocm/apt/4.1/ xenial main > /etc/apt/sources.list.d/rocm.list'
 
 # Install dependencies
 RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \
@@ -57,14 +57,7 @@ RUN pip3 install -r /doc-requirements.txt
 # Install dependencies
 ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
-# Manually ignore rocm dependencies
-RUN cget -p $PREFIX ignore \
-    RadeonOpenCompute/clang-ocl \
-    ROCm-Developer-Tools/HIP \
-    ROCmSoftwarePlatform/MIOpen \
-    ROCmSoftwarePlatform/MIOpenGEMM \
-    ROCmSoftwarePlatform/rocBLAS
-RUN cget -p $PREFIX init --cxx /opt/rocm/llvm/bin/clang++ --cc /opt/rocm/llvm/bin/clang
-RUN cget -p $PREFIX install -f dev-requirements.txt
-RUN cget -p $PREFIX install oneapi-src/oneDNN@v1.7
+
+COPY ./tools/install_prereqs.sh /
+RUN /install_prereqs.sh /usr/local / && rm /install_prereqs.sh
 

requirements.txt

 google/protobuf@v3.11.0 -DCMAKE_POSITION_INDEPENDENT_CODE=On -X subdir -Dprotobuf_BUILD_TESTS=Off
-RadeonOpenCompute/rocm-cmake@4e5e4dd --build
+RadeonOpenCompute/rocm-cmake@1cd57a0 --build
 ROCmSoftwarePlatform/rocBLAS@abd98a2b48b29326ebaef471630786a548622c06
 ROCmSoftwarePlatform/MIOpen@2.4.0
 nlohmann/json@v3.8.0

src/include/migraphx/cloneable.hpp

@@ -19,7 +19,7 @@ struct cloneable
     {
         friend Derived;
 
-        std::shared_ptr<Base> clone()
+        std::shared_ptr<Base> clone() override
         {
             return std::make_shared<Derived>(static_cast<const Derived&>(*this));
         }
@@ -31,7 +31,7 @@ struct cloneable
 
     struct share : Base, std::enable_shared_from_this<Base>
     {
-        std::shared_ptr<Base> clone() { return this->shared_from_this(); }
+        std::shared_ptr<Base> clone() override { return this->shared_from_this(); }
         template <typename... Args>
         share(Args&&... args) : Base(std::forward<Args>(args)...)
         {

src/include/migraphx/iterator_for.hpp

@@ -59,18 +59,24 @@ struct iterator_for_range
 
     struct iterator
     {
+        using difference_type   = std::ptrdiff_t;
+        using reference         = decltype(std::declval<base_iterator>());
+        using value_type        = std::remove_reference_t<reference>;
+        using pointer           = std::add_pointer_t<value_type>;
+        using iterator_category = std::input_iterator_tag;
         base_iterator i;
         auto operator*() const { return Selector::deref(i); }
         base_iterator operator++() { return ++i; }
+        bool operator==(const iterator& rhs) const { return i == rhs.i; }
         bool operator!=(const iterator& rhs) const { return i != rhs.i; }
     };
 
-    iterator begin()
+    iterator begin() const
     {
         assert(base != nullptr);
         return {Selector::begin(base)};
     }
-    iterator end()
+    iterator end() const
     {
         assert(base != nullptr);
         return {Selector::end(base)};

src/opt/memory_coloring_impl.hpp

@@ -40,10 +40,6 @@ struct live_interval
 {
     live_interval() : segment({invalid_offset, invalid_offset, invalid_offset, invalid_offset, 0})
     {
-        id               = invalid_offset;
-        def_point        = invalid_offset;
-        is_literal       = false;
-        is_live_on_entry = false;
     }
 
     void add_use(std::size_t use) { use_points.push_front(use); }
@@ -56,12 +52,12 @@ struct live_interval
 #endif
 
     live_range segment;
-    std::size_t id;
-    std::list<std::size_t> use_points;
-    std::size_t def_point;
-    shape result;
-    bool is_literal;
-    bool is_live_on_entry;
+    std::size_t id = invalid_offset;
+    std::list<std::size_t> use_points{};
+    std::size_t def_point = invalid_offset;
+    shape result{};
+    bool is_literal       = false;
+    bool is_live_on_entry = false;
 };
 
 using interval_ptr = live_interval*;
@@ -71,15 +67,6 @@ struct memory_coloring_impl
     memory_coloring_impl(module* p, std::string alloc_op, bool p_verify)
         : p_mod(p), allocation_op(std::move(alloc_op)), enable_verify(p_verify)
     {
-        instr2_live.clear();
-        live_ranges.clear();
-        conflict_table.clear();
-        num_of_lives       = 0;
-        max_value_number   = -1;
-        required_bytes     = 0;
-        earliest_end_point = -1;
-        latest_end_point   = -1;
-        unify_literals     = false;
     }
 
     bool allocate(interval_ptr);
@@ -155,23 +142,23 @@ struct memory_coloring_impl
     module* p_mod;
     std::unordered_map<const instruction*, interval_ptr> instr2_live;
     // universe of live intervals.
-    std::vector<live_interval> live_intervals;
+    std::vector<live_interval> live_intervals = {};
     // Map live range value number to live range.
-    std::unordered_map<int, live_range*> live_ranges;
+    std::unordered_map<int, live_range*> live_ranges = {};
     // Map live range value number to a set of conflicting live ranges' value numbers.
-    std::unordered_map<int, std::set<int>> conflict_table;
+    std::unordered_map<int, std::set<int>> conflict_table = {};
     // Priority queue for coloring.
-    std::priority_queue<interval_ptr, std::vector<interval_ptr>, ordering> alloc_queue;
+    std::priority_queue<interval_ptr, std::vector<interval_ptr>, ordering> alloc_queue{};
 
-    int num_of_lives;
-    int max_value_number;
-    std::size_t required_bytes;
+    int num_of_lives           = 0;
+    int max_value_number       = -1;
+    std::size_t required_bytes = 0;
     // The earliest program point where an live interval ends.
-    int earliest_end_point;
+    int earliest_end_point = -1;
     // The latest program point where an live interval ends.
-    int latest_end_point;
+    int latest_end_point = -1;
     // Whether to unify literals into coloring.
-    bool unify_literals;
+    bool unify_literals = false;
     std::string allocation_op{};
     bool enable_verify;
 

src/rewrite_rnn.cpp

@@ -49,6 +49,7 @@ void rewrite_rnn::apply(module& prog) const
     }
 }
 
+// NOLINTNEXTLINE(readability-function-cognitive-complexity)
 void rewrite_rnn::apply_vanilla_rnn(module& prog, instruction_ref ins) const
 {
     assert(ins->name() == "rnn");
@@ -356,6 +357,7 @@ std::vector<operation> rewrite_rnn::vanilla_rnn_actv_funcs(instruction_ref ins)
     }
 }
 
+// NOLINTNEXTLINE(readability-function-cognitive-complexity)
 void rewrite_rnn::apply_gru(module& prog, instruction_ref ins) const
 {
     assert(ins->name() == "gru");
@@ -533,6 +535,7 @@ void rewrite_rnn::apply_gru(module& prog, instruction_ref ins) const
     replace_last_hs_output(prog, ins, seq_lens, last_output, dirct);
 }
 
+// NOLINTNEXTLINE(readability-function-cognitive-complexity)
 std::vector<instruction_ref> rewrite_rnn::gru_cell(bool is_forward,
                                                    module& prog,
                                                    instruction_ref ins,
@@ -744,6 +747,7 @@ std::vector<operation> rewrite_rnn::gru_actv_funcs(instruction_ref ins) const
 }
 
 // for lstm operators
+// NOLINTNEXTLINE(readability-function-cognitive-complexity)
 void rewrite_rnn::apply_lstm(module& prog, instruction_ref ins) const
 {
     assert(ins->name() == "lstm");
@@ -1001,6 +1005,7 @@ void rewrite_rnn::apply_lstm(module& prog, instruction_ref ins) const
     replace_last_cell_output(prog, ins, seq_lens, cell_outputs, last_cell_output, dirct);
 }
 
+// NOLINTNEXTLINE(readability-function-cognitive-complexity)
 std::vector<instruction_ref> rewrite_rnn::lstm_cell(bool is_forward,
                                                     module& prog,
                                                     instruction_ref ins,

src/schedule.cpp

@@ -288,20 +288,12 @@ struct stream_info
     {
         return [=](auto f) {
             return fix<bool>([&](auto self, auto ins) {
-                for(auto i : select(ins))
-                {
+                return all_of(select(ins), [&](auto i) {
                     if(iweights.at(i) == 0)
-                    {
-                        if(not self(i))
-                            return false;
-                    }
+                        return self(i);
                     else
-                    {
-                        if(not f(this->get_stream(i)))
-                            return false;
-                    }
-                }
-                return true;
+                        return f(this->get_stream(i));
+                });
             })(start);
         };
     }

src/targets/gpu/CMakeLists.txt

@@ -263,7 +263,8 @@ message(STATUS "extractkernel: ${MIGRAPHX_EXTRACT_KERNEL}")
 include(TargetFlags)
 target_flags(HIP_COMPILER_FLAGS hip::device)
 # Remove cuda arch flags
-string(REGEX REPLACE --cuda-gpu-arch=[a-z0-9]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
+string(REGEX REPLACE "--cuda-gpu-arch=[^ \t\r\n]+" "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
+string(REGEX REPLACE "--offload-arch=[^ \t\r\n]+" "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
 message(STATUS "Hip compiler flags: ${HIP_COMPILER_FLAGS}")
 target_compile_definitions(migraphx_gpu PRIVATE 
     "-DMIGRAPHX_HIP_COMPILER=${CMAKE_CXX_COMPILER}" 

src/tf/tf_parser.cpp

@@ -258,13 +258,10 @@ create_literal(shape::type_t shape_type, const std::vector<size_t>& dims, std::v
 static bool is_valid_op(const tensorflow::NodeDef& node)
 {
     std::vector<std::string> ignored{"NoOp", "Assert"};
-    for(const auto& op : ignored)
-    {
+    return none_of(ignored, [&](const auto& op) {
         const auto& name = get_name(node);
-        if(node.op() == op or contains(name, op))
-            return false;
-    }
-    return true;
+        return node.op() == op or contains(name, op);
+    });
 }
 
 std::vector<std::string> tf_parser::find_outputs() const

src/value.cpp

@@ -22,16 +22,16 @@ struct value_base_impl : cloneable<value_base_impl>
     value_base_impl()                       = default;
     value_base_impl(const value_base_impl&) = default;
     value_base_impl& operator=(const value_base_impl&) = default;
-    virtual ~value_base_impl() {}
+    virtual ~value_base_impl() override {}
 };
 
-#define MIGRAPHX_VALUE_GENERATE_BASE_TYPE(vt, cpp_type)               \
-    struct vt##_value_holder : value_base_impl::share                 \
-    {                                                                 \
-        vt##_value_holder(cpp_type d) : data(std::move(d)) {}         \
-        virtual value::type_t get_type() { return value::vt##_type; } \
-        virtual const cpp_type* if_##vt() const { return &data; }     \
-        cpp_type data;                                                \
+#define MIGRAPHX_VALUE_GENERATE_BASE_TYPE(vt, cpp_type)                        \
+    struct vt##_value_holder : value_base_impl::share                          \
+    {                                                                          \
+        vt##_value_holder(cpp_type d) : data(std::move(d)) {}                  \
+        virtual value::type_t get_type() override { return value::vt##_type; } \
+        virtual const cpp_type* if_##vt() const override { return &data; }     \
+        cpp_type data;                                                         \
     };
 MIGRAPHX_VISIT_VALUE_TYPES(MIGRAPHX_VALUE_GENERATE_BASE_TYPE)
 
@@ -39,8 +39,8 @@ struct array_value_holder : value_base_impl::derive<array_value_holder>
 {
     array_value_holder() {}
     array_value_holder(std::vector<value> d) : data(std::move(d)) {}
-    virtual value::type_t get_type() { return value::array_type; }
-    virtual std::vector<value>* if_array() { return &data; }
+    virtual value::type_t get_type() override { return value::array_type; }
+    virtual std::vector<value>* if_array() override { return &data; }
     std::vector<value> data;
 };
 
@@ -51,9 +51,9 @@ struct object_value_holder : value_base_impl::derive<object_value_holder>
         : data(std::move(d)), lookup(std::move(l))
     {
     }
-    virtual value::type_t get_type() { return value::object_type; }
-    virtual std::vector<value>* if_array() { return &data; }
-    virtual std::unordered_map<std::string, std::size_t>* if_object() { return &lookup; }
+    virtual value::type_t get_type() override { return value::object_type; }
+    virtual std::vector<value>* if_array() override { return &data; }
+    virtual std::unordered_map<std::string, std::size_t>* if_object() override { return &lookup; }
     std::vector<value> data;
     std::unordered_map<std::string, std::size_t> lookup;
 };

test/schedule_test.cpp

@@ -144,17 +144,15 @@ struct schedule_model_test
 
 bool check_conflicts(migraphx::module& m, migraphx::instruction_ref x, migraphx::instruction_ref y)
 {
-    for(auto ins : migraphx::iterator_for(m))
-    {
+    return migraphx::any_of(migraphx::iterator_for(m), [&](auto ins) {
         if(ins->name() != "identity")
-            continue;
+            return false;
         if(not migraphx::contains(ins->inputs(), x))
-            continue;
+            return false;
         if(not migraphx::contains(ins->inputs(), y))
-            continue;
+            return false;
         return true;
-    }
-    return false;
+    });
 }
 
 struct scheduler

tools/install_prereqs.sh

@@ -2,8 +2,8 @@
 #
 # Build MIGraphX prerequisites for docker container
 
-#install pip3 and rocm-cmake
-apt update && apt install -y python3-pip rocm-cmake
+#install pip3, rocm-cmake, rocblas and miopen
+apt update && apt install -y python3-pip rocm-cmake rocblas miopen-hip openmp-extras
 
 # install onnx package for unit tests
 pip3 install onnx==1.7.0 numpy==1.18.5 typing==3.7.4 pytest==6.0.1
@@ -11,9 +11,6 @@ pip3 install onnx==1.7.0 numpy==1.18.5 typing==3.7.4 pytest==6.0.1
 # install rbuild to build dependencies
 pip3 install https://github.com/RadeonOpenCompute/rbuild/archive/master.tar.gz
 
-# rocblas and miopen
-apt update && apt install -y rocblas miopen-hip 
-
 PREFIX=/usr/local
 REQ_FILE_DIR=""
 if [ "$#" -ge 2 ]; then
@@ -32,7 +29,7 @@ cget -p $PREFIX ignore \
     ROCmSoftwarePlatform/MIOpen \
     ROCmSoftwarePlatform/MIOpenGEMM \
     ROCmSoftwarePlatform/rocBLAS
-cget -p $PREFIX init --cxx /opt/rocm/llvm/bin/clang++
+cget -p $PREFIX init --cxx /opt/rocm/llvm/bin/clang++ --cc /opt/rocm/llvm/bin/clang
 cget -p $PREFIX install -f ${REQ_FILE_DIR}dev-requirements.txt
 cget -p $PREFIX install oneapi-src/oneDNN@v1.7