Add depthwise Deconvolution. (#187)

* add depthwise deconvolution.

* add depthwise deconvolution.

* fix some syntax error and uncessary modification

src/CMakeLists.txt

@@ -137,6 +137,7 @@ ncnn_add_layer(Permute)
 ncnn_add_layer(PriorBox)
 ncnn_add_layer(DetectionOutput)
 ncnn_add_layer(Interp)
+ncnn_add_layer(DeconvolutionDepthWise)
 
 add_library(ncnn STATIC ${ncnn_SRCS})
 

src/layer/arm/deconvolution_3x3.h

@@ -237,3 +237,131 @@ static void deconv3x3s1_neon(const Mat& bottom_blob, Mat& top_blob, const Mat& _
         }
     }
 }
+
+static void deconv3x3s2_neon(const Mat& bottom_blob, Mat& top_blob, const Mat& _kernel, const Mat& _bias)
+{
+    int w = bottom_blob.w;
+    int h = bottom_blob.h;
+    int inch = bottom_blob.c;
+
+    int outw = top_blob.w;
+    int outch = top_blob.c;
+
+    const float* kernel = _kernel;
+    const float* bias = _bias;
+
+    #pragma omp parallel for
+    for (int p=0; p<outch; p++)
+    {
+        Mat out = top_blob.channel(p);
+
+        const float bias0 = bias ? bias[p] : 0.f;
+
+        out.fill(bias0);
+
+        for (int q=0; q<inch; q++)
+        {
+            const float* img0 = bottom_blob.channel(q);
+
+            const float* kernel0 = kernel + p*inch*9 + q*9;
+
+            const float* r0 = img0;
+
+            const float* k0 = kernel0;
+            const float* k1 = kernel0 + 3;
+            const float* k2 = kernel0 + 6;
+
+#if __ARM_NEON
+            float32x4_t _k0 = vld1q_f32(k0);
+            float32x4_t _k1 = vld1q_f32(k1);
+            float32x4_t _k2 = vld1q_f32(k2);
+#endif // __ARM_NEON
+
+            for (int i = 0; i < h; i++)
+            {
+                float* outptr = out.data + outw * i*2;
+
+                float* outptr0 = outptr;
+                float* outptr1 = outptr0 + outw;
+                float* outptr2 = outptr1 + outw;
+
+                int j = 0;
+#if __ARM_NEON
+                for (; j+3 < w; j+=4)
+                {
+                    float32x4_t _v = vld1q_f32(r0);
+
+                    // out row 0
+                    float32x4_t _out00 = vmulq_lane_f32(_v, vget_low_f32(_k0), 0);   // 0,2,4,6
+                    float32x4_t _out01 = vmulq_lane_f32(_v, vget_low_f32(_k0), 1);   // 1,3,5,7
+                    float32x4_t _out02 = vmulq_lane_f32(_v, vget_high_f32(_k0), 0);   // 2,4,6,8
+
+                    float32x4x2_t _out0 = vld2q_f32(outptr0);
+                    _out0.val[0] = vaddq_f32(_out0.val[0], _out00);     // 0,2,4,6
+                    _out0.val[1] = vaddq_f32(_out0.val[1], _out01);     // 1,3,5,7
+                    vst2q_f32(outptr0, _out0);
+
+                    _out0 = vld2q_f32(outptr0 + 2);
+                    _out0.val[0] = vaddq_f32(_out0.val[0], _out02);     // 2,4,6,8
+                    vst2q_f32(outptr0 + 2, _out0);
+
+                    // out row 1
+                    float32x4_t _out10 = vmulq_lane_f32(_v, vget_low_f32(_k1), 0);    // 0,2,4,6
+                    float32x4_t _out11 = vmulq_lane_f32(_v, vget_low_f32(_k1), 1);    // 1,3,5,7
+                    float32x4_t _out12 = vmulq_lane_f32(_v, vget_high_f32(_k1), 0);    // 2,4,6,8
+
+                    float32x4x2_t _out1 = vld2q_f32(outptr1);
+                    _out1.val[0] = vaddq_f32(_out1.val[0], _out10);     // 0,2,4,6
+                    _out1.val[1] = vaddq_f32(_out1.val[1], _out11);     // 1,3,5,7
+                    vst2q_f32(outptr1, _out1);
+
+                    _out1 = vld2q_f32(outptr1 + 2);
+                    _out1.val[0] = vaddq_f32(_out1.val[0], _out12);      // 2,4,6,8
+                    vst2q_f32(outptr1 + 2, _out1);
+
+                    // out row 2
+                    float32x4_t _out20 = vmulq_lane_f32(_v, vget_low_f32(_k2), 0);     // 0,2,4,6
+                    float32x4_t _out21 = vmulq_lane_f32(_v, vget_low_f32(_k2), 1);     // 1,3,5,7
+                    float32x4_t _out22 = vmulq_lane_f32(_v, vget_high_f32(_k2), 0);     // 2,4,6,8
+
+                    float32x4x2_t _out2 = vld2q_f32(outptr2);
+                    _out2.val[0] = vaddq_f32(_out2.val[0], _out20);     // 0,2,4,6
+                    _out2.val[1] = vaddq_f32(_out2.val[1], _out21);     // 1,3,5,7
+                    vst2q_f32(outptr2, _out2);
+
+                    _out2 = vld2q_f32(outptr2 + 2);
+                    _out2.val[0] = vaddq_f32(_out2.val[0], _out22);     // 2,4,6,8
+                    vst2q_f32(outptr2 + 2, _out2);
+
+                    r0 += 4;
+                    outptr0 += 8;
+                    outptr1 += 8;
+                    outptr2 += 8;
+                }
+#endif  // __ARM_NEON
+
+                for (; j < w; j++)
+                {
+                    float val = r0[0];
+
+                    outptr0[0] += val * k0[0];
+                    outptr0[1] += val * k0[1];
+                    outptr0[2] += val * k0[2];
+
+                    outptr1[0] += val * k1[0];
+                    outptr1[1] += val * k1[1];
+                    outptr1[2] += val * k1[2];
+
+                    outptr2[0] += val * k2[0];
+                    outptr2[1] += val * k2[1];
+                    outptr2[2] += val * k2[2];
+
+                    r0++;
+                    outptr0 += 2;
+                    outptr1 += 2;
+                    outptr2 += 2;
+                }
+            }
+        }
+    }
+}

src/layer/arm/deconvolution_4x4.h

@@ -0,0 +1,339 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#if __ARM_NEON
+#include <arm_neon.h>
+#endif // __ARM_NEON
+
+static void deconv4x4s1_neon(const Mat& bottom_blob, Mat& top_blob, const Mat& _kernel, const Mat& _bias)
+{
+    int w = bottom_blob.w;
+    int h = bottom_blob.h;
+    int inch = bottom_blob.c;
+
+    int outw = top_blob.w;
+    int outch = top_blob.c;
+
+    const float* kernel = _kernel;
+    const float* bias = _bias;
+
+    #pragma omp parallel for
+    for (int p=0; p<outch; p++)
+    {
+        Mat out = top_blob.channel(p);
+
+        const float bias0 = bias ? bias[p] : 0.f;
+
+        out.fill(bias0);
+
+        for (int q=0; q<inch; q++)
+        {
+            const float* img0 = bottom_blob.channel(q);
+
+            const float* kernel0 = kernel + p*inch*16 + q*16;
+
+            const float* r0 = img0;
+
+            const float* k0 = kernel0;
+            const float* k1 = kernel0 + 4;
+            const float* k2 = kernel0 + 8;
+            const float* k3 = kernel0 + 12;
+
+#if __ARM_NEON
+            float32x4_t _k0 = vld1q_f32(k0);
+            float32x4_t _k1 = vld1q_f32(k1);
+            float32x4_t _k2 = vld1q_f32(k2);
+            float32x4_t _k3 = vld1q_f32(k3);
+#endif // __ARM_NEON
+
+            for (int i = 0; i < h; i++)
+            {
+                float* outptr = out.data + out.w * i;
+
+                float* outptr0 = outptr;
+                float* outptr1 = outptr0 + outw;
+                float* outptr2 = outptr1 + outw;
+                float* outptr3 = outptr2 + outw;
+
+                int j = 0;
+
+#if __ARM_NEON
+                for (; j+3<w; j+=4)
+                {
+                    float32x4_t _v = vld1q_f32(r0);
+
+                    //
+                    float32x4_t _out00 = vld1q_f32(outptr0 + 0);
+                    _out00 = vmlaq_lane_f32(_out00, _v, vget_low_f32(_k0), 0);
+                    vst1q_f32(outptr0 + 0, _out00);
+
+                    float32x4_t _out01 = vld1q_f32(outptr0 + 1);
+                    _out01 = vmlaq_lane_f32(_out01, _v, vget_low_f32(_k0), 1);
+                    vst1q_f32(outptr0 + 1, _out01);
+
+                    float32x4_t _out02 = vld1q_f32(outptr0 + 2);
+                    _out02 = vmlaq_lane_f32(_out02, _v, vget_high_f32(_k0), 0);
+                    vst1q_f32(outptr0 + 2, _out02);
+
+                    float32x4_t _out03 = vld1q_f32(outptr0 + 3);
+                    _out03 = vmlaq_lane_f32(_out03, _v, vget_high_f32(_k0), 1);
+                    vst1q_f32(outptr0 + 3, _out03);
+
+                    //
+                    float32x4_t _out10 = vld1q_f32(outptr1 + 0);
+                    _out10 = vmlaq_lane_f32(_out10, _v, vget_low_f32(_k1), 0);
+                    vst1q_f32(outptr1 + 0, _out10);
+
+                    float32x4_t _out11 = vld1q_f32(outptr1 + 1);
+                    _out11 = vmlaq_lane_f32(_out11, _v, vget_low_f32(_k1), 1);
+                    vst1q_f32(outptr1 + 1, _out11);
+
+                    float32x4_t _out12 = vld1q_f32(outptr1 + 2);
+                    _out12 = vmlaq_lane_f32(_out12, _v, vget_high_f32(_k1), 0);
+                    vst1q_f32(outptr1 + 2, _out12);
+
+                    float32x4_t _out13 = vld1q_f32(outptr1 + 3);
+                    _out13 = vmlaq_lane_f32(_out13, _v, vget_high_f32(_k1), 1);
+                    vst1q_f32(outptr1 + 3, _out13);
+
+                    //
+                    float32x4_t _out20 = vld1q_f32(outptr2 + 0);
+                    _out20 = vmlaq_lane_f32(_out20, _v, vget_low_f32(_k2), 0);
+                    vst1q_f32(outptr2 + 0, _out20);
+
+                    float32x4_t _out21 = vld1q_f32(outptr2 + 1);
+                    _out21 = vmlaq_lane_f32(_out21, _v, vget_low_f32(_k2), 1);
+                    vst1q_f32(outptr2 + 1, _out21);
+
+                    float32x4_t _out22 = vld1q_f32(outptr2 + 2);
+                    _out22 = vmlaq_lane_f32(_out22, _v, vget_high_f32(_k2), 0);
+                    vst1q_f32(outptr2 + 2, _out22);
+
+                    float32x4_t _out23 = vld1q_f32(outptr2 + 3);
+                    _out23 = vmlaq_lane_f32(_out23, _v, vget_high_f32(_k2), 1);
+                    vst1q_f32(outptr2 + 3, _out23);
+
+                    //
+                    float32x4_t _out30 = vld1q_f32(outptr3 + 0);
+                    _out30 = vmlaq_lane_f32(_out30, _v, vget_low_f32(_k3), 0);
+                    vst1q_f32(outptr3 + 0, _out30);
+
+                    float32x4_t _out31 = vld1q_f32(outptr3 + 1);
+                    _out31 = vmlaq_lane_f32(_out31, _v, vget_low_f32(_k3), 1);
+                    vst1q_f32(outptr3 + 1, _out31);
+
+                    float32x4_t _out32 = vld1q_f32(outptr3 + 2);
+                    _out32 = vmlaq_lane_f32(_out32, _v, vget_high_f32(_k3), 0);
+                    vst1q_f32(outptr3 + 2, _out32);
+
+                    float32x4_t _out33 = vld1q_f32(outptr3 + 3);
+                    _out33 = vmlaq_lane_f32(_out33, _v, vget_high_f32(_k3), 1);
+                    vst1q_f32(outptr3 + 3, _out33);
+
+                    r0 += 4;
+                    outptr0 += 4;
+                    outptr1 += 4;
+                    outptr2 += 4;
+                    outptr3 += 4;
+                }
+
+#endif  // __ARM_NEON
+
+                for (; j < w; j++)
+                {
+                    float val = r0[0];
+
+                    outptr0[0] += val * k0[0];
+                    outptr0[1] += val * k0[1];
+                    outptr0[2] += val * k0[2];
+                    outptr0[3] += val * k0[3];
+
+                    outptr1[0] += val * k1[0];
+                    outptr1[1] += val * k1[1];
+                    outptr1[2] += val * k1[2];
+                    outptr1[3] += val * k1[3];
+
+                    outptr2[0] += val * k2[0];
+                    outptr2[1] += val * k2[1];
+                    outptr2[2] += val * k2[2];
+                    outptr2[3] += val * k2[3];
+
+                    outptr3[0] += val * k3[0];
+                    outptr3[1] += val * k3[1];
+                    outptr3[2] += val * k3[2];
+                    outptr3[3] += val * k3[3];
+
+                    r0++;
+                    outptr0++;
+                    outptr1++;
+                    outptr2++;
+                    outptr3++;
+                }
+            }
+        }
+    }
+}
+
+static void deconv4x4s2_neon(const Mat& bottom_blob, Mat& top_blob, const Mat& _kernel, const Mat& _bias)
+{
+    int w = bottom_blob.w;
+    int h = bottom_blob.h;
+    int inch = bottom_blob.c;
+
+    int outw = top_blob.w;
+    int outch = top_blob.c;
+
+    const float* kernel = _kernel;
+    const float* bias = _bias;
+
+    #pragma omp parallel for
+    for (int p=0; p<outch; p++)
+    {
+        Mat out = top_blob.channel(p);
+
+        const float bias0 = bias ? bias[p] : 0.f;
+
+        out.fill(bias0);
+
+        for (int q=0; q<inch; q++)
+        {
+            const float* img0 = bottom_blob.channel(q);
+
+            const float* kernel0 = kernel + p*inch*16 + q*16;
+
+            const float* r0 = img0;
+
+            const float* k0 = kernel0;
+            const float* k1 = kernel0 + 4;
+            const float* k2 = kernel0 + 8;
+            const float* k3 = kernel0 + 12;
+
+#if __ARM_NEON
+            float32x4_t _k0 = vld1q_f32(k0);
+            float32x4_t _k1 = vld1q_f32(k1);
+            float32x4_t _k2 = vld1q_f32(k2);
+            float32x4_t _k3 = vld1q_f32(k3);
+#endif // __ARM_NEON
+
+            for (int i = 0; i < h; i++)
+            {
+                float* outptr = out.data + out.w * i*2;
+
+                float* outptr0 = outptr;
+                float* outptr1 = outptr0 + outw;
+                float* outptr2 = outptr1 + outw;
+                float* outptr3 = outptr2 + outw;
+
+                int j = 0;
+#if __ARM_NEON
+                for (; j+3<w; j+=4)
+                {
+                    float32x4_t _v = vld1q_f32(r0);
+
+                    // row 0
+                    float32x4x2_t _out0 = vld2q_f32(outptr0);
+                    // 0,2,4,6
+                    _out0.val[0] = vmlaq_lane_f32(_out0.val[0], _v, vget_low_f32(_k0), 0);
+                    // 1,3,5,7
+                    _out0.val[1] = vmlaq_lane_f32(_out0.val[1], _v, vget_low_f32(_k0), 1);
+                    vst2q_f32(outptr0, _out0);
+
+                    _out0 = vld2q_f32(outptr0 + 2);
+                    // 2,4,6,8
+                    _out0.val[0] = vmlaq_lane_f32(_out0.val[0], _v, vget_high_f32(_k0), 0);
+                    // 3,5,7,9
+                    _out0.val[1] = vmlaq_lane_f32(_out0.val[1], _v, vget_high_f32(_k0), 1);
+                    vst2q_f32(outptr0 + 2, _out0);
+
+                    // row 1
+                    float32x4x2_t _out1 = vld2q_f32(outptr1);
+                    // 0,2,4,6
+                    _out1.val[0] = vmlaq_lane_f32(_out1.val[0], _v, vget_low_f32(_k1), 0);
+                    // 1,3,5,7
+                    _out1.val[1] = vmlaq_lane_f32(_out1.val[1], _v, vget_low_f32(_k1), 1);
+                    vst2q_f32(outptr1, _out1);
+
+                    _out1 = vld2q_f32(outptr1 + 2);
+                    // 2,4,6,8
+                    _out1.val[0] = vmlaq_lane_f32(_out1.val[0], _v, vget_high_f32(_k1), 0);
+                    // 3,5,7,9
+                    _out1.val[1] = vmlaq_lane_f32(_out1.val[1], _v, vget_high_f32(_k1), 1);
+                    vst2q_f32(outptr1 + 2, _out1);
+
+                    // row 2
+                    float32x4x2_t _out2 = vld2q_f32(outptr2);
+                    _out2.val[0] = vmlaq_lane_f32(_out2.val[0], _v, vget_low_f32(_k2), 0);
+                    _out2.val[1] = vmlaq_lane_f32(_out2.val[1], _v, vget_low_f32(_k2), 1);
+                    vst2q_f32(outptr2, _out2);
+
+                    _out2 = vld2q_f32(outptr1 + 2);
+                    _out2.val[0] = vmlaq_lane_f32(_out2.val[0], _v, vget_high_f32(_k2), 0);
+                    _out2.val[1] = vmlaq_lane_f32(_out2.val[1], _v, vget_high_f32(_k2), 1);
+                    vst2q_f32(outptr2 + 2, _out2);
+
+                    // row 3
+                    float32x4x2_t _out3 = vld2q_f32(outptr3);
+                    _out3.val[0] = vmlaq_lane_f32(_out3.val[0], _v, vget_low_f32(_k3), 0);
+                    _out3.val[1] = vmlaq_lane_f32(_out3.val[1], _v, vget_low_f32(_k3), 1);
+                    vst2q_f32(outptr3, _out3);
+
+                    _out3 = vld2q_f32(outptr3 + 2);
+                    _out3.val[0] = vmlaq_lane_f32(_out3.val[0], _v, vget_high_f32(_k3), 0);
+                    _out3.val[1] = vmlaq_lane_f32(_out3.val[1], _v, vget_high_f32(_k3), 1);
+                    vst2q_f32(outptr3 + 2, _out3);
+
+                    r0 += 4;
+                    outptr0 += 8;
+                    outptr1 += 8;
+                    outptr2 += 8;
+                    outptr3 += 8;
+                }
+
+#endif // __ARM_NEON
+
+                for (; j < w; j++)
+                {
+                    float val = r0[0];
+
+                    outptr0[0] += val * k0[0];
+                    outptr0[1] += val * k0[1];
+                    outptr0[2] += val * k0[2];
+                    outptr0[3] += val * k0[3];
+
+                    outptr1[0] += val * k1[0];
+                    outptr1[1] += val * k1[1];
+                    outptr1[2] += val * k1[2];
+                    outptr1[3] += val * k1[3];
+
+                    outptr2[0] += val * k2[0];
+                    outptr2[1] += val * k2[1];
+                    outptr2[2] += val * k2[2];
+                    outptr2[3] += val * k2[3];
+
+                    outptr3[0] += val * k3[0];
+                    outptr3[1] += val * k3[1];
+                    outptr3[2] += val * k3[2];
+                    outptr3[3] += val * k3[3];
+
+                    r0++;
+                    outptr0 += 2;
+                    outptr1 += 2;
+                    outptr2 += 2;
+                    outptr3 += 2;
+                }
+
+            }
+        }
+    }
+}

src/layer/arm/deconvolution_arm.cpp

@@ -16,24 +16,44 @@
 
 namespace ncnn {
 
+#include "deconvolution_4x4.h"
 #include "deconvolution_3x3.h"
 
 DEFINE_LAYER_CREATOR(Deconvolution_arm)
 
 int Deconvolution_arm::forward(const Mat& bottom_blob, Mat& top_blob) const
 {
-    if (kernel_size != 3 || stride != 1 || dilation != 1)
+    // deconvolv with NxN kernel
+    // value = value + bias
+
+    if ((kernel_size != 3 && kernel_size != 4) || stride > 2 || dilation != 1)
     {
         return Deconvolution::forward(bottom_blob, top_blob);
     }
 
     typedef void (*deconv_func)(const Mat&, Mat&, const Mat&, const Mat&);
 
-    deconv_func deconv = deconv3x3s1_neon;
+    // kernel_size x stride
+    deconv_func deconv_func_table[2][2] =
+    {
+        {
+            deconv3x3s1_neon,
+            deconv3x3s2_neon
+        },  // kernel_size = 3
+        {
+            deconv4x4s1_neon,
+            deconv4x4s2_neon
+        }   // kernel_size = 4
+    };
+
+    deconv_func deconv = deconv_func_table[kernel_size-3][stride-1];
+    if (!deconv)
+    {
+        return Deconvolution::forward(bottom_blob, top_blob);
+    }
 
     int w = bottom_blob.w;
     int h = bottom_blob.h;
-    int channels = bottom_blob.c;
 
     int outw = (w - 1) * stride + kernel_size;
     int outh = (h - 1) * stride + kernel_size;

src/layer/arm/deconvolutiondepthwise_arm.cpp

@@ -0,0 +1,133 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#include "deconvolutiondepthwise_arm.h"
+
+#ifdef _OPENMP
+#include <omp.h>
+#endif
+
+namespace ncnn {
+
+#include "deconvolution_3x3.h"
+#include "deconvolution_4x4.h"
+
+DEFINE_LAYER_CREATOR(DeconvolutionDepthWise_arm)
+
+int DeconvolutionDepthWise_arm::forward(const Mat& bottom_blob, Mat& top_blob) const
+{
+    // convolv with NxN kernel
+    // value = value + bias
+
+    if ((kernel_size != 3 && kernel_size != 4) || stride > 2 || dilation != 1)
+    {
+        return DeconvolutionDepthWise::forward(bottom_blob, top_blob);
+    }
+
+    typedef void (*deconv_func)(const Mat&, Mat&, const Mat&, const Mat&);
+
+    // kernel_size x stride
+    deconv_func deconv_func_table[2][2] =
+    {
+        {
+            deconv3x3s1_neon,
+            deconv3x3s2_neon
+        },  // kernel_size = 3
+        {
+            deconv4x4s1_neon,
+            deconv4x4s2_neon
+        }   // kernel_size = 4
+    };
+
+    deconv_func deconv = deconv_func_table[kernel_size-3][stride-1];
+    if (!deconv)
+    {
+        return DeconvolutionDepthWise::forward(bottom_blob, top_blob);
+    }
+
+    int w = bottom_blob.w;
+    int h = bottom_blob.h;
+    int channels = bottom_blob.c;
+
+    const int kernel_extent = dilation * (kernel_size - 1) + 1;
+
+    int outw = (w - 1) * stride + kernel_extent;
+    int outh = (h - 1) * stride + kernel_extent;
+
+    Mat top_blob_bordered(outw, outh, num_output);
+    if (top_blob_bordered.empty())
+        return -100;
+
+    const int maxk = kernel_size * kernel_size;
+
+    // depth-wise
+    if (channels == group && group == num_output)
+    {
+#ifdef _OPENMP
+        int nested_current = omp_get_nested();
+        omp_set_nested(0);
+#endif
+
+        #pragma omp parallel for
+        for (int g=0; g<group; g++)
+        {
+            Mat top_blob_bordered_g = top_blob_bordered.channel(g);
+            Mat bottom_blob_g = bottom_blob.channel(g);
+            Mat weight_data_g(maxk, (float*)(weight_data + maxk * g));
+
+            Mat bias_data_g;
+            if (bias_term)
+                bias_data_g = Mat(1, (float*)(bias_data + g));
+
+            deconv(bottom_blob_g, top_blob_bordered_g, weight_data_g, bias_data_g);
+        }
+
+#ifdef _OPENMP
+        omp_set_nested(nested_current);
+#endif
+        return 0;
+    }
+
+    const int channels_g = channels / group;
+    const int num_output_g = num_output / group;
+
+    for (int g=0; g<group; g++)
+    {
+        Mat top_blob_bordered_g(outw, outh, num_output_g, top_blob_bordered.channel(num_output_g * g));
+        Mat bottom_blob_g(w, h, channels_g, bottom_blob.channel(channels_g * g).data);
+        Mat weight_data_g(maxk * channels_g * num_output_g, (float*)(weight_data + maxk * channels_g * num_output_g * g));
+        Mat bias_data_g;
+        if (bias_term)
+            bias_data_g = Mat(num_output_g, (float*)(bias_data + num_output_g * g));
+
+        deconv(bottom_blob_g, top_blob_bordered_g, weight_data_g, bias_data_g);
+    }
+
+    top_blob = top_blob_bordered;
+
+    if (pad > 0)
+    {
+        copy_cut_border(top_blob_bordered, top_blob, pad, pad, pad, pad);
+        if (top_blob.empty())
+            return -100;
+
+        outw = top_blob.w;
+        outh = top_blob.h;
+    }
+
+    return 0;
+
+}
+
+}// namespace ncnn

src/layer/arm/deconvolutiondepthwise_arm.h

@@ -0,0 +1,30 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#ifndef LAYER_DECONVOLUTIONDEPTHWISE_ARM_H
+#define LAYER_DECONVOLUTIONDEPTHWISE_ARM_H
+
+#include "deconvolutiondepthwise.h"
+
+namespace ncnn {
+
+class DeconvolutionDepthWise_arm : public DeconvolutionDepthWise
+{
+public:
+    virtual int forward(const Mat& bottom_blob, Mat& top_blob) const;
+};
+
+} // namespace ncnn
+
+#endif // LAYER_DECONVOLUTIONDEPTHWISE_ARM_H

src/layer/deconvolutiondepthwise.cpp

@@ -0,0 +1,181 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#include "deconvolutiondepthwise.h"
+
+namespace ncnn {
+
+DEFINE_LAYER_CREATOR(DeconvolutionDepthWise)
+
+DeconvolutionDepthWise::DeconvolutionDepthWise()
+{
+    one_blob_only = true;
+    support_inplace = false;
+}
+
+DeconvolutionDepthWise::~DeconvolutionDepthWise()
+{
+}
+
+int DeconvolutionDepthWise::load_param(const ParamDict& pd)
+{
+    Deconvolution::load_param(pd);
+
+    group = pd.get(7, 1);
+
+    return 0;
+}
+
+int DeconvolutionDepthWise::forward(const Mat& bottom_blob, Mat& top_blob) const
+{
+    if (group == 1)
+    {
+        return Deconvolution::forward(bottom_blob, top_blob);
+    }
+
+    // deconvolv with NxN kernel
+    // value = value + bias
+
+    int w = bottom_blob.w;
+    int h = bottom_blob.h;
+    int channels = bottom_blob.c;
+
+    if (channels % group != 0 || num_output % group != 0)
+    {
+        // reject invalid group
+        return -100;
+    }
+
+    const int kernel_extent = dilation * (kernel_size - 1) + 1;
+
+    int outw = (w - 1) * stride + kernel_extent;
+    int outh = (h - 1) * stride + kernel_extent;
+
+    Mat top_blob_bordered;
+    top_blob_bordered.create(outw, outh, num_output);
+    if (top_blob_bordered.empty())
+        return -100;
+
+    const int maxk = kernel_size * kernel_size;
+
+    // kernel offsets
+    std::vector<int> _space_ofs(maxk);
+    int* space_ofs = &_space_ofs[0];
+    {
+        int p1 = 0;
+        int p2 = 0;
+        int gap = outw * dilation - kernel_size * dilation;;
+        for (int i = 0; i < kernel_size; i++)
+        {
+            for (int j = 0; j < kernel_size; j++)
+            {
+                space_ofs[p1] = p2;
+                p1++;
+                p2 += dilation;
+            }
+            p2 += gap;
+        }
+    }
+
+    // depth-wise
+    if (channels == group && group == num_output)
+    {
+        #pragma omp parallel for
+        for (int g=0; g<group; g++)
+        {
+            const float* inptr = bottom_blob.channel(g);
+            const float* kptr = weight_data + maxk * g;
+            Mat m = top_blob_bordered.channel(g);
+
+            const float bias = bias_term ? bias_data.data[g] : 0.f;
+
+            m.fill(bias);
+
+            for (int i = 0; i < h; i++)
+            {
+                for (int j = 0; j < w; j++)
+                {
+                    float* outptr = m.data + outw * i*stride + j*stride;
+
+                    for (int k = 0; k < maxk; k++)
+                    {
+                        float val = inptr[i*w + j];
+                        float w = kptr[k];
+                        outptr[ space_ofs[k] ] += val * w;
+                    }
+                }
+            }
+        }
+    }
+    else
+    {
+        // num_output
+        const int channels_g = channels / group;
+        const int num_output_g = num_output / group;
+
+        #pragma omp parallel for
+        for (int g = 0; g < group; g++)
+        {
+            const float* weight_data_ptr = weight_data + maxk * channels_g * num_output_g * g;
+            for (int p = 0; p < num_output_g; p++)
+            {
+                Mat out = top_blob_bordered.channel(g * num_output_g + p);
+
+                const float bias = bias_term ? bias_data.data[g * num_output_g + p] : 0.f;
+
+                out.fill(bias);
+
+                for (int i = 0; i < h; i++)
+                {
+                    for (int j = 0; j < w; j++)
+                    {
+                        float* outptr = out.data + out.w * i*stride + j*stride;
+
+                        const float* kptr = weight_data_ptr + maxk * channels_g * p;
+
+                        // channels_g
+                        for (int q = 0; q < channels_g; q++)
+                        {
+                            const Mat m = bottom_blob.channel(channels_g * g + q);
+                            float val = *(m.data + w * i + j);
+
+                            for (int k = 0; k < maxk; k++)
+                            {
+                                outptr[ space_ofs[k] ] += val * kptr[k];
+                            }
+
+                            kptr += maxk;
+                        }
+                    }
+                }
+            }
+        }
+    }
+
+    top_blob = top_blob_bordered;
+
+    if (pad > 0)
+    {
+        copy_cut_border(top_blob_bordered, top_blob, pad, pad, pad, pad);
+        if (top_blob.empty())
+            return -100;
+
+        outw = top_blob.w;
+        outh = top_blob.h;
+    }
+
+    return 0;
+}
+
+} // namespace ncnn

src/layer/deconvolutiondepthwise.h

@@ -0,0 +1,39 @@
+// Tencent is pleased to support the open source community by making ncnn available.
+//
+// Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
+//
+// Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
+// in compliance with the License. You may obtain a copy of the License at
+//
+// https://opensource.org/licenses/BSD-3-Clause
+//
+// Unless required by applicable law or agreed to in writing, software distributed
+// under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
+// CONDITIONS OF ANY KIND, either express or implied. See the License for the
+// specific language governing permissions and limitations under the License.
+
+#ifndef LAYER_DECONVOLUTIONDEPTHWISE_H
+#define LAYER_DECONVOLUTIONDEPTHWISE_H
+
+#include "layer.h"
+#include "deconvolution.h"
+
+namespace ncnn {
+
+class DeconvolutionDepthWise : public Deconvolution
+{
+public:
+    DeconvolutionDepthWise();
+    virtual ~DeconvolutionDepthWise();
+
+    virtual int load_param(const ParamDict& pd);
+
+    virtual int forward(const Mat& bottom_blob, Mat& top_blob) const;
+
+public:
+    int group;
+};
+
+} // namespace ncnn
+
+#endif // LAYER_DECONVOLUTIONDEPTHWISE_H

tools/caffe/caffe2ncnn.cpp

@@ -332,6 +332,14 @@ int main(int argc, char** argv)
             else
                 fprintf(pp, "%-16s", "Convolution");
         }
+        else if (layer.type() == "Deconvolution")
+        {
+            const caffe::ConvolutionParameter& convolution_param = layer.convolution_param();
+            if (convolution_param.group() != 1)
+                fprintf(pp, "%-16s", "DeconvolutionDepthWise");
+            else
+                fprintf(pp, "%-16s", "Deconvolution");
+        }
         else if (layer.type() == "Python")
         {
             const caffe::PythonParameter& python_param = layer.python_param();
@@ -542,6 +550,11 @@ int main(int argc, char** argv)
             fprintf(pp, " 5=%d", convolution_param.bias_term());
             fprintf(pp, " 6=%d", weight_blob.data_size());
 
+            if (convolution_param.group() != 1)
+            {
+                fprintf(pp, " 7=%d", convolution_param.group());
+            }
+
             int quantized_weight = 0;
             fwrite(&quantized_weight, sizeof(int), 1, bp);