新智元导读】今天给大家安利一个宝藏仓库miemiedetection , 该仓库集合了PPYOLO、PPYOLOv2、PPYOLOE三个算法pytorch实现三合一,其中的PPYOLOv2和PPYOLO算法刚刚支持了导出ncnn。
众所周知,PPYOLO和PPYOLOv2的导出部署非常困难,因为它们使用了可变形卷积、MatrixNMS等对部署不太友好的算子。
而作者在ncnn中实现了可变形卷积DCNv2、CoordConcat、PPYOLO Decode MatrixNMS等自定义层,使得使用ncnn部署PPYOLO和PPYOLOv2成为了可能。其中的可变形卷积层也已经被合入ncnn官方仓库。
在ncnn中对图片预处理时,先将图片从BGR格式转成RGB格式,然后用cv2.INTER_CUBIC方式将图片插值成640x640的大小,再使用相同的均值和标准差对图片进行归一化。以上全部与原版PPYOLOv2一样,从而确保了C++端和python端输入神经网络的图片张量是完全一样的。
最后,ncnn的输出与miemiedetection的输出对比如下图所示:
其中,右边是miemiedetection的输出,为ppyolov2_r50vd_365e.pth这个模型预测的结果。在miemiedetection根目录下输入以下内容即可得到。
python tools/demo.py image -f exps/ppyolo/ppyolov2_r50vd_365e.py -c ppyolov2_r50vd_365e.pth --path assets/000000013659.jpg --conf 0.15 --tsize 640 --save_result --device gpu
左边则是ncnn相同的模型ppyolov2_r50vd_365e的结果,ncnn的运算结果与pytorch有细微差别,影响不大。
pytorch直接转ncnn
读了一部分ncnn的源码,确保对 *.bin 和 *.param 文件充分了解之后,封装了1个工具ncnn_utils,源码位于miemiedetection的mmdet/models/ncnn_utils.py,它支持写一次前向传播就能导出ncnn使用的 *.bin 和 *.param 文件,你只需给每个pytorch层增加1个export_ncnn()方法,export_ncnn()方法几乎只要照抄farward()方法就能把模型导出到ncnn。
以下是ncnn_utils工具的使用示例:
是不是很牛x?你只要照着farward()方法写,在export_ncnn()方法里用ncnn_utils的api写一次前向传播就能把pytorch模型导出到ncnn。在这个示例中,我展示了如何将resnet中使用的ConvNormLayer层导出到ncnn,ConvNormLayer层里包含了卷积层、bn层、激活层(当self.dcn_v2==False),或者是卷积层、可变形卷积层、bn层、激活层(当self.dcn_v2==True)。为了提升ncnn的推理速度,我将卷积层(可变形卷积层)和bn层合并,另外,当激活函数是relu、leakyrelu、clip、sigmoid、mish、hardswish这些时,还可以将激活层合并到卷积层当中,这样就将3个层合并成了1个层,大大提高推理速度。
可变形卷积
卷积层可以视为可变形卷积在offset==0,mask==1时的特例。一个形状为[in_c, h, w]的特征图inputs,经过普通卷积层(卷积核形状是[num_output, in_c, kernel_h, kernel_w],w方向的步长、相邻卷积采样点的距离、卷积步长、左填充、右填充分别是kernel_w、dilation_w、stride_w、pad_left、pad_right,h方向的步长、相邻卷积采样点的距离、卷积步长、上填充、下填充分别是kernel_h、dilation_h、stride_h、pad_top、pad_bottom)后,得到的特征图形状是[num_output, out_h, out_w],其中out_h = (h + pad_top + pad_bottom - dilation_h * (kernel_h - 1) + 1) / stride_h + 1,out_w = (w + pad_left + pad_right - dilation_w * (kernel_w - 1) + 1) / stride_w + 1。一个形状为[in_c, h, w]的特征图inputs,经过可变形卷积层(卷积核形状是[num_output, in_c, kernel_h, kernel_w],w方向的步长、相邻卷积采样点的距离、卷积步长、左填充、右填充分别是kernel_w、dilation_w、stride_w、pad_left、pad_right,h方向的步长、相邻卷积采样点的距离、卷积步长、上填充、下填充分别是kernel_h、dilation_h、stride_h、pad_top、pad_bottom)后,得到的特征图形状也是[num_output, out_h, out_w],其中out_h = (h + pad_top + pad_bottom - dilation_h * (kernel_h - 1) + 1) / stride_h + 1,out_w = (w + pad_left + pad_right - dilation_w * (kernel_w - 1) + 1) / stride_w + 1。但不同的是在可变形卷积层之前,inputs需要经过一个普通卷积层,获得可变形卷积需要的offset和mask,offset和mask的形状分别是[kernel_h * kernel_w * 2, out_h, out_w]、[kernel_h * kernel_w, out_h, out_w]。为什么是这个形状呢?我们知道,inputs经过卷积层,卷积窗是不是滑动了out_h * out_w次?是的,因为每一行卷积窗滑动了out_w次,每一列卷积窗滑动了out_h次,所以总共滑动了out_h * out_w次。此外,卷积采样点是不是有kernel_h * kernel_w个?是的,offset表示的是卷积窗停留在每一个位置的时候,每个卷积采样点的偏移(有y、x两个坐标),所以offset的形状是[kernel_h * kernel_w * 2, out_h, out_w]。但是,offset是浮点数,你怎么取原图inputs里的像素?双线性插值!对采样点的x、y坐标分别进行上取整和下取整,得到最近的4个采样点的坐标,然后将4个采样点的像素进行双线性插值,得到所求的像素val。mask是0到1之间的值(进入可变形卷积层之前会经过sigmoid层),表示的是每个val的重要程度,所以它的形状是[kernel_h * kernel_w, out_h, out_w]。offset和mask会和inputs一起进入可变形卷积层参与后续计算。「talk is cheap, show me the code」,我们来看一下ncnn中可变形卷积的代码!
...
#include "deformableconv2d.h" #include "fused_activation.h" namespace ncnn { DeformableConv2D::DeformableConv2D() { one_blob_only = false; support_inplace = false; } int DeformableConv2D::load_param(const ParamDict& pd) { num_output = pd.get(0, 0); kernel_w = pd.get(1, 0); kernel_h = pd.get(11, kernel_w); dilation_w = pd.get(2, 1); dilation_h = pd.get(12, dilation_w); stride_w = pd.get(3, 1); stride_h = pd.get(13, stride_w); pad_left = pd.get(4, 0); pad_right = pd.get(15, pad_left); pad_top = pd.get(14, pad_left); pad_bottom = pd.get(16, pad_top); bias_term = pd.get(5, 0); weight_data_size = pd.get(6, 0); activation_type = pd.get(9, 0); activation_params = pd.get(10, Mat()); return 0; } int DeformableConv2D::load_model(const ModelBin& mb) { weight_data = mb.load(weight_data_size, 0); if (weight_data.empty()) return -100; if (bias_term) { bias_data = mb.load(num_output, 1); if (bias_data.empty()) return -100; } return 0; } int DeformableConv2D::forward(const std::vector<Mat>& bottom_blobs, std::vector<Mat>& top_blobs, const Option& opt) const { const Mat& bottom_blob = bottom_blobs[0]; const Mat& offset = bottom_blobs[1]; const bool has_mask = (bottom_blobs.size() == 3); const int w = bottom_blob.w; const int h = bottom_blob.h; const int in_c = bottom_blob.c; const size_t elemsize = bottom_blob.elemsize; const int kernel_extent_w = dilation_w * (kernel_w - 1) + 1; const int kernel_extent_h = dilation_h * (kernel_h - 1) + 1; const int out_w = (w + pad_left + pad_right - kernel_extent_w) / stride_w + 1; const int out_h = (h + pad_top + pad_bottom - kernel_extent_h) / stride_h + 1; // output.shape is [num_output, out_h, out_w] (in python). Mat& output = top_blobs[0]; output.create(out_w, out_h, num_output, elemsize, opt.blob_allocator); if (output.empty()) return -100; const float* weight_ptr = weight_data; const float* bias_ptr = weight_data; if (bias_term) bias_ptr = bias_data; // deformable conv #pragma omp parallel for num_threads(opt.num_threads) for (int h_col = 0; h_col < out_h; h_col++) { for (int w_col = 0; w_col < out_w; w_col++) { int h_in = h_col * stride_h - pad_top; int w_in = w_col * stride_w - pad_left; for (int oc = 0; oc < num_output; oc++) { float sum = 0.f; if (bias_term) sum = bias_ptr[oc]; for (int i = 0; i < kernel_h; i++) { for (int j = 0; j < kernel_w; j++) { const float offset_h = offset.channel((i * kernel_w + j) * 2).row(h_col)[w_col]; const float offset_w = offset.channel((i * kernel_w + j) * 2 + 1).row(h_col)[w_col]; const float mask_ = has_mask ? bottom_blobs[2].channel(i * kernel_w + j).row(h_col)[w_col] : 1.f; const float h_im = h_in + i * dilation_h + offset_h; const float w_im = w_in + j * dilation_w + offset_w; // Bilinear const bool cond = h_im > -1 && w_im > -1 && h_im < h && w_im < w; int h_low = 0; int w_low = 0; int h_high = 0; int w_high = 0; float w1 = 0.f; float w2 = 0.f; float w3 = 0.f; float w4 = 0.f; bool v1_cond = false; bool v2_cond = false; bool v3_cond = false; bool v4_cond = false; if (cond) { h_low = floor(h_im); w_low = floor(w_im); h_high = h_low + 1; w_high = w_low + 1; float lh = h_im - h_low; float lw = w_im - w_low; float hh = 1 - lh; float hw = 1 - lw; v1_cond = (h_low >= 0 && w_low >= 0); v2_cond = (h_low >= 0 && w_high <= w - 1); v3_cond = (h_high <= h - 1 && w_low >= 0); v4_cond = (h_high <= h - 1 && w_high <= w - 1); w1 = hh * hw; w2 = hh * lw; w3 = lh * hw; w4 = lh * lw; } for (int c_im = 0; c_im < in_c; c_im++) { float val = 0.f; if (cond) { float v1 = v1_cond ? bottom_blob.channel(c_im).row(h_low)[w_low] : 0.f; float v2 = v2_cond ? bottom_blob.channel(c_im).row(h_low)[w_high] : 0.f; float v3 = v3_cond ? bottom_blob.channel(c_im).row(h_high)[w_low] : 0.f; float v4 = v4_cond ? bottom_blob.channel(c_im).row(h_high)[w_high] : 0.f; val = w1 * v1 + w2 * v2 + w3 * v3 + w4 * v4; } sum += val * mask_ * weight_ptr[((oc * in_c + c_im) * kernel_h + i) * kernel_w + j]; } } } output.channel(oc).row(h_col)[w_col] = activation_ss(sum, activation_type, activation_params); } } } return 0; } } // namespace ncnn