Nd deconv cpu (#565)

* initial progress

* formatting

* check existing tests

* formatting

* change for loop to transform

* formatting

* add tests

* formatting

* remove comment

* add more tests

* update gpu miopen calls

* formatting

* initial progress

* add cpu impl and tests

* formatting

* add NOLINT

* add 3d test

* formatting

* add more op_shape tests

* fix error msg

* fix bounds

* formatting

* fix algorithm

* formatting

* pin numpy version
Co-authored-by: mvermeulen <5479696+mvermeulen@users.noreply.github.com>

Dockerfile

@@ -59,7 +59,7 @@ ENV LC_ALL=C.UTF-8
 ENV LANG=C.UTF-8
 
 # Install cget
-RUN pip3 install cget && pip3 install numpy
+RUN pip3 install cget && pip3 install numpy==1.18.5
 
 # Install rclone
 RUN pip install https://github.com/pfultz2/rclone/archive/master.tar.gz

src/targets/cpu/lowering.cpp

@@ -26,6 +26,7 @@
 #include <migraphx/cpu/gemm.hpp>
 #include <unordered_map>
 #include <utility>
+#include <iostream>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
@@ -275,41 +276,69 @@ struct cpu_deconvolution
 
             std::fill(output.begin(), output.end(), type{0});
 
-            auto out_lens = output_shape.lens();
-            auto out_h    = out_lens[2];
-            auto out_w    = out_lens[3];
-
-            auto in   = input.get_shape().lens();
-            auto in_n = in[0];
-            auto in_c = in[1];
-            auto in_h = in[2];
-            auto in_w = in[3];
+            auto in_lens = input.get_shape().lens();
+            auto in_n    = in_lens[0];
+            auto in_c    = in_lens[1];
 
             auto wei   = weights.get_shape().lens();
             auto wei_n = wei[0];
             auto wei_c = wei[1];
-            auto wei_h = wei[2];
-            auto wei_w = wei[3];
-
-            par_dfor(in_n, wei_c)([&](std::size_t o, std::size_t k) {
-
-                dfor(in_c, in_h, in_w, wei_h, wei_w)(
-                    [&](std::size_t w, std::size_t i, std::size_t j, std::size_t x, std::size_t y) {
-                        const int start_x = i * op.stride[0] - op.padding[0];
-                        const int start_y = j * op.stride[1] - op.padding[1];
-                        const int out_x   = start_x + x * op.dilation[0];
-                        const int out_y   = start_y + y * op.dilation[1];
-
-                        const auto group_id = w / (wei_n / op.group);
-                        const auto in_ch    = group_id * wei_c + k;
-
-                        if(out_x >= 0 && out_x < out_h && out_y >= 0 && out_y < out_w)
-                        {
-                            output(o, in_ch, out_x, out_y) +=
-                                input(o, w, i, j) * weights(w, k, x, y);
-                        }
-                    });
+
+            auto out_lens = output_shape.lens();
+            auto kdims    = op.kdims();
+
+            std::vector<std::size_t> win_size{in_c};
+            std::copy(in_lens.begin() + 2, in_lens.end(), std::back_inserter(win_size));
+            std::copy(wei.begin() + 2, wei.end(), std::back_inserter(win_size));
+            shape win_shape{output_shape.type(), win_size};
+
+            par_dfor(in_n, wei_c)([&](int o, int k) {
+
+                shape_for_each(win_shape, [&](auto idx_win) {
+                    const int w = idx_win[0];
+
+                    auto input_dims_start = idx_win.begin() + 1;
+                    auto wei_dims_start   = idx_win.begin() + kdims + 1;
+
+                    std::vector<std::ptrdiff_t> win_start;
+                    for(std::size_t n = 0; n < kdims; ++n)
+                    {
+                        win_start.push_back(std::ptrdiff_t(*(input_dims_start + n) * op.stride[n]) -
+                                            std::ptrdiff_t(op.padding[n]));
+                    }
+
+                    const int group_id = w / (wei_n / op.group);
+                    const int in_ch    = group_id * wei_c + k;
+
+                    std::vector<std::ptrdiff_t> idx_out{o, in_ch};
+
+                    for(size_t n = 0; n < kdims; n++)
+                    {
+                        idx_out.push_back(win_start[n] + *(wei_dims_start + n) * op.dilation[n]);
+                    }
+
+                    std::vector<std::ptrdiff_t> idx_wei{w, k};
+                    std::copy(wei_dims_start, idx_win.end(), std::back_inserter(idx_wei));
+
+                    std::vector<std::ptrdiff_t> idx_in{o, w};
+                    std::copy(input_dims_start, wei_dims_start, std::back_inserter(idx_in));
+
+                    if(std::all_of(
+                           idx_out.begin() + 2, idx_out.end(), [&](auto ii) { return ii >= 0; }) and
+                       std::equal(idx_out.begin() + 2,
+                                  idx_out.end(),
+                                  out_lens.begin() + 2,
+                                  out_lens.end(),
+                                  std::less<std::ptrdiff_t>{}))
+                    {
+                        output(idx_out.begin(), idx_out.end()) +=
+                            input(idx_in.begin(), idx_in.end()) *
+                            weights(idx_wei.begin(), idx_wei.end());
+                    }
+                });
+
             });
+
         });
         return result;
     }
@@ -493,7 +522,7 @@ struct cpu_pad
         return migraphx::reflect(self.op, f);
     }
 
-    std::string name() const { return "cpu::contiguous"; }
+    std::string name() const { return "cpu::pad"; }
     shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
     argument compute(context&, const shape& output_shape, std::vector<argument> args) const
     {

test/cpu_ops_test.cpp

@@ -2193,6 +2193,71 @@ TEST_CASE(deconv_test)
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
+TEST_CASE(deconv_1d_test)
+{
+    migraphx::shape s{migraphx::shape::float_type, {1, 1, 3}};
+    std::vector<float> x_data{0, 0.5, 1};
+    std::vector<float> w_data{0.5, 0.5, 0.5};
+
+    std::vector<float> gold{0, 0.25, 0.75, 0.75, 0.5};
+
+    migraphx::program p;
+    auto x = p.add_literal(migraphx::literal{s, x_data});
+    auto w = p.add_literal(migraphx::literal{s, w_data});
+
+    p.add_instruction(migraphx::op::deconvolution{{0}, {1}, {1}}, x, w);
+    p.compile(migraphx::cpu::target{});
+    auto result = p.eval({}).back();
+
+    std::vector<float> results_vector;
+    result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+    EXPECT(migraphx::verify_range(results_vector, gold));
+}
+
+TEST_CASE(deconv_3d_test)
+{
+    migraphx::shape s_1{migraphx::shape::float_type, {1, 1, 1, 2, 3}};
+    migraphx::shape s_2{migraphx::shape::float_type, {1, 1, 3, 2, 3}};
+    std::vector<float> x_data{0.8471, -0.4195, -2.2749, 1.2491, 0.1722, 0.3246};
+    std::vector<float> w_data{0.6478,
+                              -0.1985,
+                              0.0633,
+                              -0.3479,
+                              2.7056,
+                              -0.1440,
+                              -1.1229,
+                              -0.7507,
+                              -1.3151,
+                              0.8884,
+                              -0.1859,
+                              -0.3407,
+                              -1.1544,
+                              -1.5893,
+                              1.6265,
+                              -1.4624,
+                              0.3812,
+                              -1.5378};
+
+    std::vector<float> gold{0.5488,  -0.4399, -1.3369, 0.4251,  -0.1439, 0.5145,  2.3015,  -0.2104,
+                            -6.1482, 0.3482,  -0.4346, 3.3197,  0.1731,  0.8533,  -0.0467, -0.9512,
+                            -0.1649, 1.7553,  2.2594,  2.9917,  -0.6500, -1.6612, -4.3680, 0.0957,
+                            0.3482,  1.1097,  -0.0792, -0.1692, -0.1190, -0.1106, -0.9779, -0.8621,
+                            4.6707,  2.9332,  -3.7001, -2.6808, -1.2476, 3.2475,  -0.4578, 4.0263,
+                            -1.8267, 0.2243,  -2.3299, -0.1411, -0.4991};
+
+    migraphx::program p;
+    auto x = p.add_literal(migraphx::literal{s_1, x_data});
+    auto w = p.add_literal(migraphx::literal{s_2, w_data});
+
+    p.add_instruction(migraphx::op::deconvolution{{0, 0, 0}, {1, 1, 1}, {1, 1, 1}}, x, w);
+    p.compile(migraphx::cpu::target{});
+    auto result = p.eval({}).back();
+
+    std::vector<float> results_vector;
+    result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+    EXPECT(migraphx::verify_range(results_vector, gold));
+}
+
 TEST_CASE(transpose_test)
 {
     migraphx::shape a_shape{migraphx::shape::float_type, {1, 2, 2, 3}};

test/gpu/ops_test.cpp

@@ -755,6 +755,62 @@ struct test_group_conv : verify_program<test_group_conv>
     }
 };
 
+struct test_deconv : verify_program<test_deconv>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto input =
+            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3}});
+        auto weights =
+            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3}});
+        p.add_instruction(migraphx::op::deconvolution{}, input, weights);
+        return p;
+    }
+};
+
+struct test_deconv_2x3 : verify_program<test_deconv_2x3>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto input =
+            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 6, 7}});
+        auto weights =
+            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {3, 4, 3, 3}});
+        p.add_instruction(migraphx::op::deconvolution{{1, 1}, {2, 3}, {1, 1}}, input, weights);
+        return p;
+    }
+};
+
+struct test_deconv_1d : verify_program<test_deconv_1d>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1, 3}});
+        auto weights =
+            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {1, 1, 3}});
+        p.add_instruction(migraphx::op::deconvolution{{0}, {1}, {1}}, input, weights);
+        return p;
+    }
+};
+
+struct test_deconv_3d : verify_program<test_deconv_3d>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto input =
+            p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3, 3}});
+        auto weights =
+            p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3, 3}});
+        p.add_instruction(
+            migraphx::op::deconvolution{{0, 0, 0}, {1, 1, 1}, {1, 1, 1}}, input, weights);
+        return p;
+    }
+};
+
 struct test_conv_relu : verify_program<test_conv_relu>
 {
     migraphx::program create_program() const

test/op_shape_test.cpp

@@ -69,6 +69,7 @@ TEST_CASE(convolution_shape)
     migraphx::shape weights{migraphx::shape::float_type, {4, 3, 3, 3}};
     expect_shape(output, migraphx::op::convolution{}, input, weights);
     throws_shape(migraphx::op::convolution{}, input);
+    throws_shape(migraphx::op::convolution{{0}, {1}, {1}}, input);
 
     migraphx::shape input2{migraphx::shape::float_type, {3, 3}};
     migraphx::shape weights2{migraphx::shape::float_type, {3, 3}};
@@ -98,6 +99,7 @@ TEST_CASE(deconvolution_shape)
     migraphx::shape weights{migraphx::shape::float_type, {4, 3, 3, 3}};
     expect_shape(output, migraphx::op::deconvolution{}, input, weights);
     throws_shape(migraphx::op::deconvolution{}, input);
+    throws_shape(migraphx::op::deconvolution{{0}, {1}, {1}}, input);
 
     migraphx::shape input_1d{migraphx::shape::float_type, {4, 4, 1}};
     migraphx::shape output_1d{migraphx::shape::float_type, {4, 3, 3}};
@@ -121,6 +123,7 @@ TEST_CASE(quant_convolution_shape)
     expect_shape(output, migraphx::op::quant_convolution{}, input, weights);
     throws_shape(migraphx::op::quant_convolution{}, input);
     throws_shape(migraphx::op::quant_convolution{{0}, {1, 1}, {1, 1}}, input, weights);
+    throws_shape(migraphx::op::quant_convolution{{0}, {1}, {1}}, input, weights);
 
     migraphx::shape input2{migraphx::shape::int32_type, {3, 3}};
     migraphx::shape weights2{migraphx::shape::float_type, {3, 3}};