Fix debug assert (#930)

* fix two asserts for debug build

* add unit test for copy parameters

* clang format

* add a unit test for reorder_dims

* change tranpose to always require perm not be empty

* clang format

* remove an unnecessary line

* fix tidy error

* fix review comments

src/include/migraphx/op/transpose.hpp

@@ -32,31 +32,23 @@ struct transpose
         auto input_lens    = input.lens();
         auto input_strides = input.strides();
         auto t             = input.type();
-        auto tuned_dims    = dims;
-        // if not perm provided, reverse the dims
-        if(tuned_dims.empty())
-        {
-            tuned_dims.resize(input_lens.size());
-            std::iota(tuned_dims.begin(), tuned_dims.end(), 0);
-            std::reverse(tuned_dims.begin(), tuned_dims.end());
-        }
 
-        if(tuned_dims.size() != input_lens.size())
+        if(dims.size() != input_lens.size())
         {
             MIGRAPHX_THROW("Permutation has wrong number of axes");
         }
-        std::vector<int64_t> axes(tuned_dims.size());
+        std::vector<int64_t> axes(dims.size());
         std::iota(axes.begin(), axes.end(), 0);
-        if(!std::is_permutation(axes.begin(), axes.end(), tuned_dims.begin()))
+        if(!std::is_permutation(axes.begin(), axes.end(), dims.begin()))
         {
-            MIGRAPHX_THROW("Invalid permutation");
+            MIGRAPHX_THROW("TRANSPOSE: Invalid permutation");
         }
         std::vector<size_t> output_lens(input_lens.size());
         std::vector<size_t> output_strides(input_lens.size());
         for(std::size_t i = 0; i < output_lens.size(); i++)
         {
-            output_lens[i]    = input_lens[tuned_dims[i]];
-            output_strides[i] = input_strides[tuned_dims[i]];
+            output_lens[i]    = input_lens[dims[i]];
+            output_strides[i] = input_strides[dims[i]];
         }
         return {t, output_lens, output_strides};
     }

src/onnx/parse_transpose.cpp

 #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 {
@@ -21,6 +22,20 @@ struct parse_transpose : op_parser<parse_transpose>
             auto&& perm_vals = info.attributes["perm"].ints();
             perm             = std::vector<int64_t>(perm_vals.begin(), perm_vals.end());
         }
+
+        // if perm is empty, use the default value
+        auto n_dim = args.front()->get_shape().lens().size();
+        if(perm.empty())
+        {
+            perm.resize(n_dim);
+            std::iota(perm.rbegin(), perm.rend(), 0);
+        }
+
+        if(perm.size() != n_dim)
+        {
+            MIGRAPHX_THROW("PARSE_TRANSPOSE: perm and input have diffferent number of dims!");
+        }
+
         return info.add_instruction(make_op("transpose", {{"permutation", perm}}), args.front());
     }
 };

src/targets/gpu/lowering.cpp

@@ -196,6 +196,10 @@ struct miopen_apply
             if(ins->name() != "@param")
                 continue;
 
+            // parameter no outputs, no need to insert copy to gpu
+            if(ins->outputs().empty())
+                continue;
+
             auto pos = std::next(ins);
             auto a   = insert_allocation(pos, ins->get_shape());
             auto c   = mod->insert_instruction(pos, make_op("hip::copy_to_gpu"), ins, a);

test/gpu/adjust_allocation.cpp

@@ -13,15 +13,16 @@
 #include <migraphx/op/tanh.hpp>
 #include <migraphx/op/transpose.hpp>
 #include <migraphx/pass_manager.hpp>
+#include <migraphx/make_op.hpp>
 #include <basic_ops.hpp>
 #include <test.hpp>
 
-void run_lowering(migraphx::program& p)
+void run_lowering(migraphx::program& p, bool offload_copy = false)
 {
     auto ctx = migraphx::gpu::context{};
     migraphx::run_passes(*p.get_main_module(),
                          {migraphx::auto_contiguous{},
-                          migraphx::gpu::lowering{&ctx, false},
+                          migraphx::gpu::lowering{&ctx, offload_copy},
                           migraphx::dead_code_elimination{},
                           migraphx::eliminate_contiguous{"gpu::contiguous"},
                           migraphx::dead_code_elimination{}});
@@ -67,4 +68,41 @@ TEST_CASE(tanh_shape)
     EXPECT(p1 == p2);
 }
 
+TEST_CASE(no_copy_dead_param)
+{
+    auto create_program = [] {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto x = mm->add_parameter("x", s);
+        mm->add_parameter("y", s);
+        auto sum = mm->add_instruction(migraphx::make_op("add"), x, x);
+        mm->add_return({sum});
+
+        return p;
+    };
+
+    auto create_gpu_program = [] {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape s{migraphx::shape::float_type, {2, 3}};
+        auto x = mm->add_parameter("x", s);
+        mm->add_parameter("y", s);
+        auto xb = mm->add_instruction(migraphx::make_op("hip::allocate", {{"shape", to_value(s)}}));
+        auto gx = mm->add_instruction(migraphx::make_op("hip::copy_to_gpu"), x, xb);
+        auto ab = mm->add_instruction(migraphx::make_op("hip::allocate", {{"shape", to_value(s)}}));
+        auto sum = mm->add_instruction(migraphx::make_op("gpu::add"), gx, gx, ab);
+        auto r   = mm->add_instruction(migraphx::make_op("hip::copy_from_gpu"), sum);
+        mm->add_return({r});
+
+        return p;
+    };
+
+    auto p1 = create_program();
+    auto p2 = create_gpu_program();
+
+    run_lowering(p1, true);
+    EXPECT(p1 == p2);
+}
+
 int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/onnx/gen_onnx.py

@@ -4150,6 +4150,35 @@ def tile_test_3x2():
             [helper.make_tensor('y', TensorProto.INT64, [2], [3, 2])])
 
 
+@onnx_test
+def transpose_default_perm_test():
+    x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [1, 5, 2, 3])
+    y = helper.make_tensor_value_info('1', TensorProto.FLOAT, [3, 2, 5, 1])
+
+    node = onnx.helper.make_node(
+        'Transpose',
+        inputs=['0'],
+        outputs=['1'],
+    )
+
+    return ([node], [x], [y])
+
+
+@onnx_test
+def transpose_invalid_perm_test():
+    x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [1, 2, 4, 3])
+    y = helper.make_tensor_value_info('1', TensorProto.FLOAT, [1, 3, 2, 2])
+
+    node = onnx.helper.make_node(
+        'Transpose',
+        perm=[0, 2, 1],
+        inputs=['0'],
+        outputs=['1'],
+    )
+
+    return ([node], [x], [y])
+
+
 @onnx_test
 def transpose_test():
     x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [1, 2, 2, 3])

test/onnx/onnx_test.cpp

@@ -3810,6 +3810,25 @@ TEST_CASE(tile_test_3x2)
     EXPECT(p == prog);
 }
 
+TEST_CASE(transpose_default_perm_test)
+{
+    migraphx::program p;
+    auto* mm   = p.get_main_module();
+    auto input = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 5, 2, 3}});
+    std::vector<int64_t> perm{3, 2, 1, 0};
+    auto r = mm->add_instruction(migraphx::make_op("transpose", {{"permutation", perm}}), input);
+    mm->add_return({r});
+
+    auto prog = migraphx::parse_onnx("transpose_default_perm_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+TEST_CASE(transpose_invalid_perm_test)
+{
+    EXPECT(test::throws([&] { migraphx::parse_onnx("transpose_invalid_perm_test.onnx"); }));
+}
+
 TEST_CASE(transpose_test)
 {
     migraphx::program p;

test/onnx/transpose_default_perm_test.onnx

+transpose_default_perm_test:j
+
+
0
1"	Transposetranspose_default_perm_testZ
+
0
+

+

+
+
+b
+
1
+

+
+
+
+
B
\ No newline at end of file

test/op_shape_test.cpp

@@ -1560,7 +1560,6 @@ TEST_CASE(transpose_shape)
     migraphx::shape output{migraphx::shape::float_type, {2, 2}, {1, 2}};
     expect_shape(input, migraphx::make_op("transpose", {{"permutation", {0, 1}}}), input);
     expect_shape(output, migraphx::make_op("transpose", {{"permutation", {1, 0}}}), input);
-    expect_shape(output, migraphx::make_op("transpose"), input);
     throws_shape(migraphx::make_op("transpose", {{"permutation", {1, 2}}}), input);
 }