Merge branch 'develop' of https://github.com/ROCmSoftwarePlatform/AMDMIGraphX into int8_miopen_call

src/targets/gpu/include/migraphx/gpu/softmax.hpp

@@ -34,7 +34,7 @@ struct miopen_softmax
         return migraphx::reflect(self.op, f);
     }
 
-    std::string name() const { return "gpu::softmax"; }
+    std::string name() const { return "miopen::softmax"; }
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;

src/tf/CMakeLists.txt

@@ -31,7 +31,7 @@ rocm_install_targets(
 
 add_executable(read_tf read_tf.cpp)
 rocm_clang_tidy_check(read_tf)
-target_link_libraries(read_tf migraphx_tf)
+target_link_libraries(read_tf migraphx_tf migraphx_cpu)
 
 if(MIGRAPHX_ENABLE_GPU)
 add_executable(verify_tf verify_tf.cpp)

src/tf/tf.cpp

@@ -317,6 +317,7 @@ struct tf_parser
             }
         }
 
+        auto l0 = args[0];
         if(contains(attributes, "padding"))
         {
             const std::string& pad_mode = attributes.at("padding").s();
@@ -326,8 +327,24 @@ struct tf_parser
                 std::vector<size_t> weight_dims = weights->get_shape().lens();
                 size_t weight_h                 = weight_dims[2];
                 size_t weight_w                 = weight_dims[3];
-                op.padding[0]                   = calculate_padding(weight_h, op.dilation[0]);
-                op.padding[1]                   = calculate_padding(weight_w, op.dilation[1]);
+
+                auto input_dims = l0->get_shape().lens();
+                size_t input_h  = input_dims[2];
+                size_t input_w  = input_dims[3];
+                std::vector<int64_t> pads(input_dims.size());
+                calculate_padding(0, pads, input_h, op.stride[0], op.dilation[0], weight_h);
+                calculate_padding(1, pads, input_w, op.stride[1], op.dilation[1], weight_w);
+
+                if(pads[0] != pads[2] || pads[1] != pads[3])
+                {
+                    std::vector<int64_t> padding = {0, 0, pads[0], pads[1], 0, 0, pads[2], pads[3]};
+                    l0 = prog.add_instruction(migraphx::op::pad{padding}, l0);
+                }
+                else
+                {
+                    op.padding[0] = pads[0];
+                    op.padding[1] = pads[1];
+                }
             }
             else if(pad_mode.find("VALID") != std::string::npos)
             {
@@ -350,7 +367,7 @@ struct tf_parser
             }
         }
 
-        return prog.add_instruction(op, {args[0], weights});
+        return prog.add_instruction(op, {l0, weights});
     }
 
     instruction_ref parse_depthwiseconv(const std::string&,
@@ -400,17 +417,35 @@ struct tf_parser
             }
         }
 
+        auto l0 = args[0];
         if(contains(attributes, "padding"))
         {
-            const std::string& pad_mode     = attributes.at("padding").s();
-            std::vector<size_t> weight_dims = weights->get_shape().lens();
-            size_t weight_h                 = weight_dims[2];
-            size_t weight_w                 = weight_dims[3];
+            const std::string& pad_mode = attributes.at("padding").s();
+
             if(pad_mode.find("SAME") != std::string::npos)
             {
-                op.padding_mode = op::padding_mode_t::same;
-                op.padding[0]   = calculate_padding(weight_h, op.dilation[0]);
-                op.padding[1]   = calculate_padding(weight_w, op.dilation[1]);
+                op.padding_mode                 = op::padding_mode_t::same;
+                std::vector<size_t> weight_dims = weights->get_shape().lens();
+                size_t weight_h                 = weight_dims[2];
+                size_t weight_w                 = weight_dims[3];
+
+                auto input_dims = l0->get_shape().lens();
+                size_t input_h  = input_dims[2];
+                size_t input_w  = input_dims[3];
+                std::vector<int64_t> pads(input_dims.size());
+                calculate_padding(0, pads, input_h, op.stride[0], op.dilation[0], weight_h);
+                calculate_padding(1, pads, input_w, op.stride[1], op.dilation[1], weight_w);
+
+                if(pads[0] != pads[2] || pads[1] != pads[3])
+                {
+                    std::vector<int64_t> padding = {0, 0, pads[0], pads[1], 0, 0, pads[2], pads[3]};
+                    l0 = prog.add_instruction(migraphx::op::pad{padding}, l0);
+                }
+                else
+                {
+                    op.padding[0] = pads[0];
+                    op.padding[1] = pads[1];
+                }
             }
             else if(pad_mode.find("VALID") != std::string::npos)
             {
@@ -432,7 +467,7 @@ struct tf_parser
         auto cweights    = prog.add_instruction(op::contiguous{}, weights);
         auto new_weights = prog.add_instruction(op::reshape{new_weights_shape}, cweights);
 
-        return prog.add_instruction(op, {args[0], new_weights});
+        return prog.add_instruction(op, {l0, new_weights});
     }
 
     instruction_ref
@@ -567,21 +602,39 @@ struct tf_parser
             op.lengths[0] = ksize[2];
             op.lengths[1] = ksize[3];
         }
+
+        auto l0 = args[0];
         if(contains(attributes, "padding"))
         {
             const std::string& pad_mode = attributes.at("padding").s();
             if(pad_mode.find("SAME") != std::string::npos)
             {
                 op.padding_mode = op::padding_mode_t::same;
-                op.padding[0]   = calculate_padding(op.lengths[0], 1);
-                op.padding[1]   = calculate_padding(op.lengths[1], 1);
+                auto input_dims = l0->get_shape().lens();
+                size_t input_h  = input_dims[2];
+                size_t input_w  = input_dims[3];
+                std::vector<int64_t> pads(input_dims.size());
+                calculate_padding(0, pads, input_h, op.stride[0], 1, op.lengths[0]);
+                calculate_padding(1, pads, input_w, op.stride[1], 1, op.lengths[1]);
+
+                if(pads[0] != pads[2] || pads[1] != pads[3])
+                {
+                    std::vector<int64_t> padding = {0, 0, pads[0], pads[1], 0, 0, pads[2], pads[3]};
+                    l0                           = prog.add_instruction(
+                        migraphx::op::pad{padding, std::numeric_limits<float>::lowest()}, l0);
+                }
+                else
+                {
+                    op.padding[0] = pads[0];
+                    op.padding[1] = pads[1];
+                }
             }
             else if(pad_mode.find("VALID") != std::string::npos)
             {
                 op.padding_mode = op::padding_mode_t::valid;
             }
         }
-        return prog.add_instruction(op, args[0]);
+        return prog.add_instruction(op, l0);
     }
 
     instruction_ref

test/CMakeLists.txt

@@ -119,7 +119,7 @@ foreach(ONNX_TEST ${ONNX_TESTS})
     set(TEST_NAME test_${BASE_NAME})
     add_executable(${TEST_NAME} ${TES_ONNX_DIR}/${ONNX_TEST})
     rocm_clang_tidy_check(${TEST_NAME})
-    target_link_libraries(${TEST_NAME} migraphx_onnx)
+    target_link_libraries(${TEST_NAME} migraphx_onnx migraphx_cpu)
     target_include_directories(${TEST_NAME} PUBLIC include)
     add_test(NAME ${TEST_NAME} COMMAND $<TARGET_FILE:${TEST_NAME}> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/onnx) 
     add_dependencies(tests ${TEST_NAME})
@@ -129,7 +129,7 @@ endforeach()
 # tf test
 add_executable(test_tf tf/tf_test.cpp)
 rocm_clang_tidy_check(test_tf)
-target_link_libraries(test_tf migraphx_tf)
+target_link_libraries(test_tf migraphx_tf migraphx_cpu)
 target_include_directories(test_tf PUBLIC include)
 add_test(NAME test_tf COMMAND $<TARGET_FILE:test_tf> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tf) 
 add_dependencies(tests test_tf)

test/eliminate_pad_test.cpp

@@ -83,23 +83,4 @@ TEST_CASE(rewrite_test_asymmetric)
         p.begin(), p.end(), [](const migraphx::instruction& ins) { return ins.name() == "pad"; }));
 }
 
-TEST_CASE(rewrite_test_same_padding)
-{
-    migraphx::program p;
-    size_t img_dim[2] = {2, 2};
-    size_t channels   = 1;
-    std::vector<int32_t> input(channels * img_dim[0] * img_dim[1]);
-    std::iota(input.begin(), input.end(), 0);
-
-    migraphx::shape s_img{migraphx::shape::int32_type, {1, channels, img_dim[0], img_dim[1]}};
-    auto l_img      = p.add_literal(migraphx::literal{s_img, input});
-    auto padded_img = p.add_instruction(migraphx::op::pad{{0, 0, 1, 1, 0, 0, 1, 1}}, l_img);
-
-    create_conv(padded_img, channels, p, migraphx::op::padding_mode_t::same);
-
-    p.compile(eliminate_pad_target{});
-    EXPECT(std::any_of(
-        p.begin(), p.end(), [](const migraphx::instruction& ins) { return ins.name() == "pad"; }));
-}
-
 int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/gpu/miopen.cpp

@@ -1646,6 +1646,22 @@ struct test_pad : verify_program<test_pad>
     }
 };
 
+struct test_pad_int8 : verify_program<test_pad_int8>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        std::vector<int8_t> data0 = {0, 1, 2, 3};
+        migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
+        auto l0 = p.add_literal(migraphx::literal{s0, data0});
+        migraphx::op::pad op{};
+        op.value = std::numeric_limits<int8_t>::lowest();
+        op.pads  = {0, 0, 1, 1};
+        p.add_instruction(op, l0);
+        return p;
+    }
+};
+
 struct test_pooling_autopad : verify_program<test_pooling_autopad>
 {
     migraphx::program create_program() const