基于darknet框架实现DepthwiseConvolutional
具体详见:https://github.com/ChenYingpeng/darknet-mobilenet ,里面包含mobilenet v1 的配置文件以及模型,在Imagenet上测试的top1:72.03%, top5:90.514%.
一、添加depthwise_convolutional_layer
1 在 /src/parser.c 中添加以下代码:
1)添加头文件
#include "utils.h"
++ #include "depthwise_convolutional_layer.h" //added by chen
2)在 string_to_layer_type函数中, 添加网络层类型解析:
if (strcmp(type, "[upsample]")==0) return UPSAMPLE;
++ if (strcmp(type, "[depthwise_convolutional]") == 0) return DEPTHWISE_CONVOLUTIONAL; //added by chen
return BLANK;
3)在 parse_network_cfg 函数中,添加网络层解析后进行构建:
if(lt == CONVOLUTIONAL){
l = parse_convolutional(options, params);
}
++ else if (lt == DEPTHWISE_CONVOLUTIONAL) {
++ l = parse_depthwise_convolutional(options, params); //added by chen
++ }
else if(lt == DECONVOLUTIONAL){
l = parse_deconvolutional(options, params);
}
其中 parse_depthwise_convolutional 需要自己编写,以下是其代码:
//added by chen
depthwise_convolutional_layer parse_depthwise_convolutional(list *options, size_params params)
{
int size = option_find_int(options, "size", 1);
int stride = option_find_int(options, "stride", 1);
int pad = option_find_int_quiet(options, "pad", 0);
int padding = option_find_int_quiet(options, "padding", 0);
if (pad) padding = size / 2;
char *activation_s = option_find_str(options, "activation", "logistic");
ACTIVATION activation = get_activation(activation_s);
int batch, h, w, c;
h = params.h;
w = params.w;
c = params.c;
batch = params.batch;
if (!(h && w && c)) error("Layer before convolutional layer must output image.");
int batch_normalize = option_find_int_quiet(options, "batch_normalize", 0);
depthwise_convolutional_layer layer = make_depthwise_convolutional_layer(batch, h, w, c, size, stride, padding, activation, batch_normalize);
layer.flipped = option_find_int_quiet(options, "flipped", 0);
layer.dot = option_find_float_quiet(options, "dot", 0);
return layer;
}
4) 在 load_weights_upto函数中,添加网络层读取参数函数;
if (l.dontload) continue;
++ if (l.type == DEPTHWISE_CONVOLUTIONAL) {
++ load_depthwise_convolutional_weights(l, fp);//added by chen
++ }
其中 load_depthwise_convolutional_weights 函数需要自己编写,实现如下:
//added by chen
void load_depthwise_convolutional_weights(layer l, FILE *fp)
{
int num = l.n*l.size*l.size;
fread(l.biases, sizeof(float), l.n, fp);
if (l.batch_normalize && (!l.dontloadscales)) {
fread(l.scales, sizeof(float), l.n, fp);
fread(l.rolling_mean, sizeof(float), l.n, fp);
fread(l.rolling_variance, sizeof(float), l.n, fp);
if (0) {
int i;
for (i = 0; i < l.n; ++i) {
printf("%g, ", l.rolling_mean[i]);
}
printf("\n");
for (i = 0; i < l.n; ++i) {
printf("%g, ", l.rolling_variance[i]);
}
printf("\n");
}
if (0) {
fill_cpu(l.n, 0, l.rolling_mean, 1);
fill_cpu(l.n, 0, l.rolling_variance, 1);
}
}
fread(l.weights, sizeof(float), num, fp);
if (l.flipped) {
//transpose_matrix(l.weights, l.c*l.size*l.size, l.n);//hjimce
}
#ifdef GPU
if (gpu_index >= 0) {
push_depthwise_convolutional_layer(l);
}
#endif
}
5)在 save_weights_upto 函数中,添加网络层保存参数函数;
if (l.dontsave) continue;
++ if (l.type == DEPTHWISE_CONVOLUTIONAL) {
++ save_depthwise_convolutional_weights(l, fp); //added by chen
++ }
其中 save_depthwise_convolutional_weights 函数需要自己编写,实现如下:
//added by chen
void save_depthwise_convolutional_weights(layer l, FILE *fp)
{
#ifdef GPU
if (gpu_index >= 0) {
pull_depthwise_convolutional_layer(l);
}
#endif
int num = l.n*l.size*l.size;
fwrite(l.biases, sizeof(float), l.n, fp);
if (l.batch_normalize) {
fwrite(l.scales, sizeof(float), l.n, fp);
fwrite(l.rolling_mean, sizeof(float), l.n, fp);
fwrite(l.rolling_variance, sizeof(float), l.n, fp);
}
fwrite(l.weights, sizeof(float), num, fp);
}
2 在 /src/network.c 中,添加以下代码;
1)添加头文件;
#include "data.h"
++ #include "depthwise_convolutional_layer.h" // added by chen
2) 添加网络层;
if(l.type == CONVOLUTIONAL){
resize_convolutional_layer(&l, w, h);
}
++ else if(l.type == DEPTHWISE_CONVOLUTIONAL){
++ resize_depthwise_convolutional_layer(&l, w, h); //added by chen
++ }
3 在 /include/darknet.h 中,添加网络层枚举类型;
在darknet.h文件中枚举类型LAYER_TYPE,添加网络层枚举类型,代码如下;
BLANK,
++ DEPTHWISE_CONVOLUTIONAL //added by chen
4 往/src/目录下添加DepthwiseConvolutional实现代码depthwise_convolutional_layer.h depthwise_convolutional_layer.c depthwise_convolutional_kernels.cu;
/src/depthwise_convolutional_layer.h
#ifndef DEPTHWISE_CONVOLUTIONAL_LAYER_H
#define DEPTHWISE_CONVOLUTIONAL_LAYER_H
#include "cuda.h"
#include "image.h"
#include "activations.h"
#include "layer.h"
#include "network.h"
typedef layer depthwise_convolutional_layer;
#ifdef GPU
void forward_depthwise_convolutional_layer_gpu(depthwise_convolutional_layer layer, network net);
void backward_depthwise_convolutional_layer_gpu(depthwise_convolutional_layer layer, network net);
void update_depthwise_convolutional_layer_gpu(depthwise_convolutional_layer layer, update_args a);
void push_depthwise_convolutional_layer(depthwise_convolutional_layer layer);
void pull_depthwise_convolutional_layer(depthwise_convolutional_layer layer);
void add_bias_gpu(float *output, float *biases, int batch, int n, int size);
void backward_bias_gpu(float *bias_updates, float *delta, int batch, int n, int size);
void adam_update_gpu(float *w, float *d, float *m, float *v, float B1, float B2, float eps, float decay, float rate, int n, int batch, int t);
#ifdef CUDNN
void cudnn_depthwise_convolutional_setup(layer *l);
#endif
#endif
depthwise_convolutional_layer make_depthwise_convolutional_layer(int batch, int h, int w, int c, int size, int stride, int padding, ACTIVATION activation, int batch_normalize);
void resize_depthwise_convolutional_layer(depthwise_convolutional_layer *layer, int w, int h);
void forward_depthwise_convolutional_layer(const depthwise_convolutional_layer layer, network net);
void update_depthwise_convolutional_layer(depthwise_convolutional_layer layer, update_args a);
void denormalize_depthwise_convolutional_layer(depthwise_convolutional_layer l);
void backward_depthwise_convolutional_layer(depthwise_convolutional_layer layer, network net);
void add_bias(float *output, float *biases, int batch, int n, int size);
void backward_bias(float *bias_updates, float *delta, int batch, int n, int size);
int depthwise_convolutional_out_height(depthwise_convolutional_layer layer);
int depthwise_convolutional_out_width(depthwise_convolutional_layer layer);
#endif
/src/depthwise_convolutional_layer.c
#include "depthwise_convolutional_layer.h"
#include "utils.h"
#include "batchnorm_layer.h"
#include "im2col.h"
#include "col2im.h"
#include "blas.h"
#include "gemm.h"
#include <stdio.h>
#include <time.h>
int depthwise_convolutional_out_height(depthwise_convolutional_layer l)
{
return (l.h + 2*l.pad - l.size) / l.stride + 1;
}
int depthwise_convolutional_out_width(depthwise_convolutional_layer l)
{
return (l.w + 2*l.pad - l.size) / l.stride + 1;
}
//ᅵᅵʱᅵᅵᅵݿՌᅵᅵÐ?
static size_t get_workspace_size(layer l){
#ifdef CUDNN
if(gpu_index >= 0){
size_t most = 0;
size_t s = 0;
cudnnGetConvolutionForwardWorkspaceSize(cudnn_handle(),
l.srcTensorDesc,
l.weightDesc,
l.convDesc,
l.dstTensorDesc,
l.fw_algo,
&s);
if (s > most) most = s;
cudnnGetConvolutionBackwardFilterWorkspaceSize(cudnn_handle(),
l.srcTensorDesc,
l.ddstTensorDesc,
l.convDesc,
l.dweightDesc,
l.bf_algo,
&s);
if (s > most) most = s;
cudnnGetConvolutionBackwardDataWorkspaceSize(cudnn_handle(),
l.weightDesc,
l.ddstTensorDesc,
l.convDesc,
l.dsrcTensorDesc,
l.bd_algo,
&s);
if (s > most) most = s;
return