Merge branch 'develop' into jit-vector-reduce

CMakeLists.txt

@@ -203,6 +203,8 @@ rocm_enable_cppcheck(
         useSmartPointer:*make_shared_array.hpp
         constParameter:*src/targets/gpu/*.cpp
         constParameter:*src/targets/gpu/*.hpp
+        # Suppress mlir_conv.cpp since this file will be deleted
+        *:*src/targets/gpu/mlir_conv.cpp
     FORCE
     INCONCLUSIVE
     RULE_FILE

dev-requirements.txt

@@ -2,6 +2,6 @@ pfultz2/rocm-recipes
 facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cmake
 ccache@v4.1
 pcre,pfultz2/pcre@8.45 -H sha256:d6f7182602a775a7d500a0cedca6449af0400c6493951513046d17615ed0bf11
-danmar/cppcheck@2.6 -DHAVE_RULES=1
+danmar/cppcheck@2.8 -DHAVE_RULES=1
 RadeonOpenCompute/rocm-cmake@1ebf7e7bc61bb5e949c171562b421264065230a7 --build
 -f requirements.txt

src/api/include/migraphx/migraphx.hpp

@@ -39,10 +39,7 @@ template <class T, class F, class... Ts>
 T* make(F f, Ts&&... xs)
 {
     T* result = nullptr;
-    // cppcheck-suppress redundantInitialization
-    // cppcheck-suppress redundantAssignment
-    // cppcheck-suppress unreadVariable
-    auto e = f(&result, std::forward<Ts>(xs)...);
+    auto e    = f(&result, std::forward<Ts>(xs)...);
     if(e != migraphx_status_success)
         throw std::runtime_error("Failed to call function");
     return result;
@@ -51,9 +48,6 @@ T* make(F f, Ts&&... xs)
 template <class F, class... Ts>
 void call(F f, Ts&&... xs)
 {
-    // cppcheck-suppress redundantInitialization
-    // cppcheck-suppress redundantAssignment
-    // cppcheck-suppress unreadVariable
     auto e = f(std::forward<Ts>(xs)...);
     if(e != migraphx_status_success)
         throw std::runtime_error("Failed to call function");
@@ -340,7 +334,6 @@ struct interface_base : Base
     template <class T, class Setter, class F>
     void set_auto_fp(Setter setter, F f)
     {
-        // cppcheck-suppress constParameter
         return set_fp<T>(setter, [=](T& obj, auto out, auto... xs) {
             auto_invoke(f, out, obj, auto_convert_param(rank<2>{}, xs)...);
         });

src/argument.cpp

@@ -29,7 +29,6 @@ void argument::assign_buffer(std::function<char*()> d)
     // Collect all shapes
     std::unordered_map<std::size_t, shape> shapes;
     {
-        // cppcheck-suppress variableScope
         std::size_t i = 0;
         fix([&](auto self, auto ss) {
             if(ss.sub_shapes().empty())
@@ -60,7 +59,6 @@ void argument::assign_buffer(std::function<char*()> d)
     }
     assert(offset == s.bytes());
 
-    // cppcheck-suppress variableScope
     std::size_t i = 0;
     m_data        = fix<data_t>([&](auto self, auto ss) {
         data_t result;

src/include/migraphx/raw_data.hpp

@@ -207,8 +207,7 @@ auto visit_all_pack(const shape& s, V1&& v1)
 template <class T, class... Ts>
 auto visit_all(T&& x, Ts&&... xs)
 {
-    auto&& s = x.get_shape();
-    // cppcheck-suppress redundantInitialization
+    auto&& s                                   = x.get_shape();
     std::initializer_list<shape::type_t> types = {xs.get_shape().type()...};
     if(!std::all_of(types.begin(), types.end(), [&](shape::type_t t) { return t == s.type(); }))
         MIGRAPHX_THROW("Types must be the same");

src/include/migraphx/serialize.hpp

@@ -50,7 +50,6 @@ auto to_value_impl(rank<2>, const T& x) -> decltype(x.begin(), x.end(), value{})
     value result = value::array{};
     for(auto&& y : x)
     {
-        auto e = to_value(y);
         result.insert(to_value(y));
     }
     return result;

src/include/migraphx/tensor_view.hpp

@@ -120,10 +120,8 @@ struct tensor_view
         return m_data[m_shape.index(this->size() - 1)];
     }
 
-    // cppcheck-suppress functionConst
     iterator begin() { return {0, {this}}; }
 
-    // cppcheck-suppress functionConst
     iterator end() { return {this->size(), {this}}; }
 
     const_iterator begin() const { return {0, {this}}; }

src/include/migraphx/verify.hpp

@@ -168,7 +168,6 @@ bool verify_range(const R1& r1, const R2& r2, double tolerance = 80, double* out
 {
     double threshold = std::numeric_limits<range_value<R1>>::epsilon() * tolerance;
     auto error       = rms_range(r1, r2);
-    // cppcheck-suppress uninitvar
     if(out_error != nullptr)
         *out_error = error;
     return error <= threshold;

src/module.cpp

@@ -737,7 +737,6 @@ std::unordered_map<instruction_ref, std::string>
 module::print_cpp(std::ostream& os, std::unordered_map<instruction_ref, std::string> names) const
 {
     os << "migraphx::module p;" << std::endl;
-    // cppcheck-suppress variableScope
     unsigned long seed = 0;
     names              = this->print(
         [&](auto ins, auto ins_names) {

src/onnx/parse_pooling.cpp

@@ -128,7 +128,7 @@ struct parse_pooling : op_parser<parse_pooling>
             std::fill_n(values["stride"].begin(), kdims, 1);
         }
         // used to calculate the supposed output shape
-        std::vector<int64_t> orig_padding(paddings.begin(), paddings.end());
+        std::vector<int64_t> orig_padding = paddings;
 
         std::vector<int64_t> slice_start;
         std::vector<int64_t> slice_end;

src/onnx/parse_squeeze.cpp

@@ -30,11 +30,11 @@ struct parse_squeeze : op_parser<parse_squeeze>
                           std::vector<instruction_ref> args) const
     {
         auto op = parser.load(opd.op_name, info);
-        std::vector<int64_t> axes;
         if(args.size() == 2)
         {
             auto arg_axes = args.at(1)->eval();
             check_arg_empty(arg_axes, "PARSE_" + opd.op_name + ": cannot handle variable axes!");
+            std::vector<int64_t> axes;
             arg_axes.visit([&](auto s) { axes.assign(s.begin(), s.end()); });
             op = assign_axes(op, axes);
         }

src/process.cpp

@@ -20,7 +20,6 @@ int exec(const std::string& cmd, const std::function<void(const char*)>& std_out
     int ec = 0;
     if(enabled(MIGRAPHX_TRACE_CMD_EXECUTE{}))
         std::cout << cmd << std::endl;
-    std::array<char, 128> buffer;
     auto closer = [&](FILE* stream) {
         auto status = pclose(stream);
         ec          = WIFEXITED(status) ? 0 : WEXITSTATUS(status); // NOLINT
@@ -30,6 +29,7 @@ int exec(const std::string& cmd, const std::function<void(const char*)>& std_out
         std::unique_ptr<FILE, decltype(closer)> pipe(popen(cmd.c_str(), "r"), closer); // NOLINT
         if(!pipe)
             MIGRAPHX_THROW("popen() failed: " + cmd);
+        std::array<char, 128> buffer;
         while(fgets(buffer.data(), buffer.size(), pipe.get()) != nullptr)
             std_out(buffer.data());
     }

src/targets/cpu/copy.cpp

@@ -20,7 +20,6 @@ struct cpu_copy : reduce_dims_base, auto_register_op<cpu_copy>
         return inputs.at(1);
     }
     argument
-    // cppcheck-suppress constParameter
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
     {
         argument result = get_arg(args, args.size() - 1);

src/targets/cpu/gather.cpp

@@ -26,7 +26,6 @@ struct cpu_gather : auto_register_op<cpu_gather>
     }
 
     argument
-    // cppcheck-suppress constParameter
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
     {
         std::size_t nelements = output_shape.elements();

src/targets/cpu/include/migraphx/cpu/pointwise.hpp

@@ -323,7 +323,6 @@ struct cpu_unary : reduce_dims_base, auto_register_op<cpu_unary<Op>>
         return {s.type(), s.lens()};
     }
     argument
-    // cppcheck-suppress constParameter
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
     {
         argument result = get_arg(args, args.size() - 1);
@@ -362,7 +361,6 @@ struct cpu_binary : reduce_dims_base, auto_register_op<cpu_binary<Op>>
     }
 
     argument
-    // cppcheck-suppress constParameter
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
     {
         argument result = get_arg(args, args.size() - 1);

src/targets/gpu/compile_hip.cpp

@@ -134,7 +134,6 @@ struct hiprtc_program
         std::vector<char> buffer(n);
         MIGRAPHX_HIPRTC(hiprtcGetProgramLog(prog.get(), buffer.data()));
         assert(buffer.back() == 0);
-        // cppcheck-suppress returnDanglingLifetime
         return {buffer.begin(), buffer.end() - 1};
     }
 

src/targets/gpu/fuse_ops.cpp

@@ -681,7 +681,7 @@ struct miopen_fusion
 struct miopen_conv_bias
 {
     op::convolution op;
-    fusion f          = {};
+    fusion fp         = {};
     fusion::op_t conv = {};
     fusion::op_t bias = {};
 
@@ -705,19 +705,19 @@ struct miopen_conv_bias
         float beta  = 0;
         miopenSetOpArgsConvForward(fargs.get(), conv, &alpha, &beta, args[1].implicit());
         miopenSetOpArgsBiasForward(fargs.get(), bias, &alpha, &beta, args[3].implicit());
-        return f.execute(ctx, fargs, args[0], args[4]);
+        return fp.execute(ctx, fargs, args[0], args[4]);
     }
 
     void finalize(context& ctx, const shape&, const std::vector<shape>& inputs)
     {
-        f    = fusion(inputs[0]);
-        conv = f.create_conv(op, inputs[1]);
-        bias = f.create_bias(inputs[3]);
-        if(not f.compile(ctx))
+        fp   = fusion(inputs[0]);
+        conv = fp.create_conv(op, inputs[1]);
+        bias = fp.create_bias(inputs[3]);
+        if(not fp.compile(ctx))
             MIGRAPHX_THROW("Failed to compile fusion plan");
     }
 
-    shape get_workspace(context& ctx) { return f.get_workspace(ctx); }
+    shape get_workspace(context& ctx) { return fp.get_workspace(ctx); }
     std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
     {
         return shapes.size() - 1;
@@ -728,7 +728,7 @@ MIGRAPHX_REGISTER_OP(miopen_conv_bias)
 struct miopen_conv_bias_relu
 {
     op::convolution op;
-    fusion f          = {};
+    fusion fp         = {};
     fusion::op_t conv = {};
     fusion::op_t bias = {};
     fusion::op_t relu = {};
@@ -754,18 +754,18 @@ struct miopen_conv_bias_relu
         miopenSetOpArgsConvForward(fargs.get(), conv, &alpha, &beta, args[1].implicit());
         miopenSetOpArgsBiasForward(fargs.get(), bias, &alpha, &beta, args[3].implicit());
         miopenSetOpArgsActivForward(fargs.get(), relu, &alpha, &beta, 0, 0, 0);
-        return f.execute(ctx, fargs, args[0], args[4]);
+        return fp.execute(ctx, fargs, args[0], args[4]);
     }
     void finalize(context& ctx, const shape&, const std::vector<shape>& inputs)
     {
-        f    = fusion(inputs[0]);
-        conv = f.create_conv(op, inputs[1]);
-        bias = f.create_bias(inputs[3]);
-        relu = f.create_relu();
-        f.compile(ctx);
+        fp   = fusion(inputs[0]);
+        conv = fp.create_conv(op, inputs[1]);
+        bias = fp.create_bias(inputs[3]);
+        relu = fp.create_relu();
+        fp.compile(ctx);
     }
 
-    shape get_workspace(context& ctx) { return f.get_workspace(ctx); }
+    shape get_workspace(context& ctx) { return fp.get_workspace(ctx); }
     std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
     {
         return shapes.size() - 1;
@@ -875,7 +875,6 @@ struct find_conv_pointwise
         {
             if(i.name()[0] == '@')
                 continue;
-            auto inputs = to_shapes(i.inputs());
             op.ops.push_back({{i.get_operator()}});
         }
         std::vector<instruction_ref> inputs = {input_ins, weights_ins, bias_ins, alloc_ins};

src/targets/gpu/pack_int8_args.cpp

@@ -22,10 +22,10 @@ static instruction_ref pad_ins(module& m, instruction_ref ins, int offset)
     auto pad_k                     = (k + 3) / 4 * 4;
     auto pad_lens                  = lens;
     pad_lens[lens.size() + offset] = pad_k;
-    std::vector<int64_t> pad_dims(lens.size() * 2, 0);
-    auto ret_ins = ins;
+    auto ret_ins                   = ins;
     if(pad_k != k)
     {
+        std::vector<int64_t> pad_dims(lens.size() * 2, 0);
         pad_dims[lens.size() + offset] = pad_k - k;
         shape ps{s.type(), pad_lens};
         auto ins_out =

test/onnx/onnx_test.cpp

@@ -3831,7 +3831,6 @@ TEST_CASE(reshape_non_standard_test)
     migraphx::program p;
     auto* mm = p.get_main_module();
     migraphx::op::reshape op;
-    std::vector<int64_t> reshape_dims{4, 3, 2};
     migraphx::shape s{migraphx::shape::float_type, {2, 3, 4}};
     auto x = mm->add_parameter("x", s);
     auto tran_x =

test/ref_rnn_ops_test.cpp

@@ -1173,12 +1173,6 @@ TEST_CASE(gru_forward_args)
         0.3852,  -0.1170, -0.2937, 0.2979,  -0.1357, 0.4257,  0.3884,  -0.2916, 0.1071,  0.0934,
         0.3645,  -0.4310, -0.3480, 0.0702,  -0.1558};
 
-    migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
-    std::vector<float> bias_data{
-        0.0560,  0.0310, -0.1669, -0.0781, 0.1793, -0.1758, 0.3173,  -0.1650, -0.3732, 0.2946,
-        -0.0912, 0.3118, 0.1391,  0.2755,  0.2695, -0.1059, -0.2357, 0.3629,  -0.2534, -0.0494,
-        0.0556,  0.0881, -0.2592, -0.2213, 0.2310, -0.4044, 0.1801,  0.1438,  0.3108,  -0.3607};
-
     migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
     std::vector<float> input{-0.8432,
                              -0.9887,
@@ -1199,9 +1193,6 @@ TEST_CASE(gru_forward_args)
                              -1.0536,
                              -0.2529};
 
-    migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
-    std::vector<float> ih_data{
-        -0.0468, 0.5691, -0.0882, 0.8340, 0.1483, -0.3902, -0.5348, 0.4178, 1.0175, 0.9212};
     float clip = 0.0f;
 
     // 3 args
@@ -1242,6 +1233,11 @@ TEST_CASE(gru_forward_args)
 
     // 4 args (bias is used)
     {
+        std::vector<float> bias_data{
+            0.0560,  0.0310, -0.1669, -0.0781, 0.1793, -0.1758, 0.3173,  -0.1650, -0.3732, 0.2946,
+            -0.0912, 0.3118, 0.1391,  0.2755,  0.2695, -0.1059, -0.2357, 0.3629,  -0.2534, -0.0494,
+            0.0556,  0.0881, -0.2592, -0.2213, 0.2310, -0.4044, 0.1801,  0.1438,  0.3108,  -0.3607};
+        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
         migraphx::program p;
         auto* mm  = p.get_main_module();
         auto seq  = mm->add_literal(migraphx::literal{in_shape, input});
@@ -1280,6 +1276,9 @@ TEST_CASE(gru_forward_args)
 
     // 4 args (ih is used)
     {
+        std::vector<float> ih_data{
+            -0.0468, 0.5691, -0.0882, 0.8340, 0.1483, -0.3902, -0.5348, 0.4178, 1.0175, 0.9212};
+        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
         migraphx::program p;
         auto* mm = p.get_main_module();
         auto seq = mm->add_literal(migraphx::literal{in_shape, input});
@@ -2210,15 +2209,6 @@ TEST_CASE(gru_bidirectional_args)
         0.4101,  0.2641,  -0.4110, -0.1681, 0.3582,  -0.2089, 0.0852,  0.0963,  0.3866,  0.1955,
         -0.2174, 0.1996,  -0.2252, 0.1748,  0.1833,  -0.3155, 0.2567,  -0.4387, 0.3402,  0.0599};
 
-    migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
-    std::vector<float> bias_data{
-        -0.1582, -0.0826, 0.4008,  0.0118,  0.2511,  0.1900,  -0.2838, 0.2549,  -0.2484, 0.2363,
-        -0.4083, -0.0295, -0.1161, 0.1211,  0.2509,  -0.1414, -0.2628, -0.2992, 0.1517,  0.1817,
-        -0.2783, 0.3183,  -0.1629, -0.3108, -0.3418, 0.0411,  0.2203,  0.2187,  -0.2990, -0.0416,
-        0.0209,  -0.1024, 0.4443,  -0.4420, -0.0330, -0.3591, -0.2990, 0.2167,  0.1395,  0.2317,
-        0.1318,  0.1909,  -0.3615, 0.1953,  -0.2582, -0.2217, 0.3723,  0.1458,  0.2630,  -0.0377,
-        0.1754,  0.0800,  -0.3964, -0.3247, 0.4219,  -0.0900, 0.3553,  0.2614,  -0.1298, -0.1124};
-
     migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
     std::vector<float> input{-0.8432,
                              -0.9887,
@@ -2239,11 +2229,6 @@ TEST_CASE(gru_bidirectional_args)
                              -1.0536,
                              -0.2529};
 
-    migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
-    std::vector<float> ih_data{-0.0468, 0.5691,  -0.0882, 0.8340,  0.1483, -0.3902, -0.5348,
-                               0.4178,  1.0175,  0.9212,  -0.0468, 0.5691, -0.0882, 0.8340,
-                               0.1483,  -0.3902, -0.5348, 0.4178,  1.0175, 0.9212};
-
     float clip = 0.0f;
 
     // 3 args
@@ -2288,6 +2273,15 @@ TEST_CASE(gru_bidirectional_args)
 
     // 4 args (bias is used)
     {
+        std::vector<float> bias_data{
+            -0.1582, -0.0826, 0.4008,  0.0118,  0.2511,  0.1900,  -0.2838, 0.2549,  -0.2484,
+            0.2363,  -0.4083, -0.0295, -0.1161, 0.1211,  0.2509,  -0.1414, -0.2628, -0.2992,
+            0.1517,  0.1817,  -0.2783, 0.3183,  -0.1629, -0.3108, -0.3418, 0.0411,  0.2203,
+            0.2187,  -0.2990, -0.0416, 0.0209,  -0.1024, 0.4443,  -0.4420, -0.0330, -0.3591,
+            -0.2990, 0.2167,  0.1395,  0.2317,  0.1318,  0.1909,  -0.3615, 0.1953,  -0.2582,
+            -0.2217, 0.3723,  0.1458,  0.2630,  -0.0377, 0.1754,  0.0800,  -0.3964, -0.3247,
+            0.4219,  -0.0900, 0.3553,  0.2614,  -0.1298, -0.1124};
+        migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
         migraphx::program p;
         auto* mm  = p.get_main_module();
         auto seq  = mm->add_literal(migraphx::literal{in_shape, input});
@@ -2330,6 +2324,10 @@ TEST_CASE(gru_bidirectional_args)
 
     // 4 args (ih is used)
     {
+        std::vector<float> ih_data{-0.0468, 0.5691,  -0.0882, 0.8340,  0.1483, -0.3902, -0.5348,
+                                   0.4178,  1.0175,  0.9212,  -0.0468, 0.5691, -0.0882, 0.8340,
+                                   0.1483,  -0.3902, -0.5348, 0.4178,  1.0175, 0.9212};
+        migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
         migraphx::program p;
         auto* mm = p.get_main_module();
         auto seq = mm->add_literal(migraphx::literal{in_shape, input});
@@ -4186,7 +4184,6 @@ TEST_CASE(lstm_bidirectional_var_seq_lens)
                                 -0.83699064, 0.49162736, -0.8271,     -0.5683,     0.4562,
                                 -1.2545,     1.2729,     -0.4082,     -0.4392,     -0.9406,
                                 0.7794,      1.8194,     -0.5811,     0.2166};
-    std::vector<int> sl_data{1, 2, 3};
 
     float clip = 0.0f;
     migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
@@ -4196,10 +4193,11 @@ TEST_CASE(lstm_bidirectional_var_seq_lens)
     migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
     migraphx::shape ic_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
     migraphx::shape pph_shape{migraphx::shape::float_type, {num_dirct, 3 * hidden_size}};
-    migraphx::shape sl_shape{migraphx::shape::int32_type, {batch_size}};
 
     // concatenation of hidden states as program output
     {
+        std::vector<int> sl_data{1, 2, 3};
+        migraphx::shape sl_shape{migraphx::shape::int32_type, {batch_size}};
         migraphx::program p;
         auto* mm    = p.get_main_module();
         auto seq    = mm->add_literal(migraphx::literal{in_shape, input_data});