【模型推理】談談 darknet yolo 的 route 算子

歡迎關注我的公衆号 [極智視界]，擷取我的更多筆記分享

O_o

>_<

o_O

O_o

~_~

o_O

本文聊一聊 darknet yolo 網絡中的 route 算子。

yolo 是目标檢測算法落地常用的網絡，具有速度快、精度高的優點，相信很多同學都熟悉，route 層在 yolov1、yolov2、yolov3、yolov4 中均有出現，yolov4 tiny 中的 route 又有了一些新特性，現在的它既能做類似 concatenate 的拼接，也能做類似 slice 的切割，算子功能十分豐富。這裡把 route 拿出來聊一聊，主要說一下它的功能、與 darknet 源碼實作。

文章目錄

- 1、route 化作 concatenate
- 2、route 化作 slice
- 3、route 化作恒等映射
- 4、darknet 源碼解讀

1、route 化作 concatenate

上面說過 route 能化作 concatenate，來完成拼接的操作，來看看 route 是怎麼樣幻化成一縷 convatenate 的，下面一張圖就可以看明白，看右側 route 參數欄可知，route 層的輸入為 4 和 5，即接兩個大卷積輸入，次元分别為 160 x 160 x 32、160 x 160 x 32，這裡 route 的作用就是把這兩個大卷積的輸出做拼接，可以看到經 route 作用後的輸出次元為 160 x 160 x 64，64 就是 32 + 32 來的。

【模型推理】談談 darknet yolo 的 route 算子

2、route 化作 slice

route 除了能化作 concatenate，還能化作 slice 來完成切割的操作，讓我們來看看 route 是怎麼樣幻化成一縷 slice 的。同樣來看下面這張圖，這時候的 route 隻有一個輸入，即 2 号大卷積，然後你再仔細觀察會發現還有 groups 和 group_id 兩個奇怪的參數，這兩個參數就是拿來做切割的參數。groups 的意思是把輸入切成 groups 份，group_id 的意思是拿第 group_id 索引的資料傳出給下一層，是以 yolo 的 route 切割其實和 slice 又不太一樣，route 切割後是要了一份丢了一份 (可能是防止過拟合)，而 slice 切割後往往是兩份都要，這個坑大家需要注意一下。

【模型推理】談談 darknet yolo 的 route 算子

3、route 化作恒等映射

這個其實是我猜的，先來說下什麼是恒等映射，恒等映射是 resnet 殘差結構中的一個概念，如下圖，恒等映射的意思就是把輸入 x 原模原樣的通過一個暢通無阻的高速路給傳遞過去。為什麼要做這個操作呢，有了恒等映射，網絡再深我也能通過這個小道把資訊給你傳遞下去，且 x 的導數為 1，也可防止回傳時梯度消失，關鍵還不增加額外計算量。

【模型推理】談談 darknet yolo 的 route 算子

回到我們的 route，來看一下下面這張圖，可以看到框起來這兩個 route 其實啥也沒做，就是原封不動的将 20 x 20 x 512 次元的資料傳遞下去，這是不是和恒等映射異曲同工。

【模型推理】談談 darknet yolo 的 route 算子

4、darknet 源碼解讀

先看下 route_layer 的頭：

#ifndef ROUTE_LAYER_H
#define ROUTE_LAYER_H
#include "network.h"
#include "layer.h"

typedef layer route_layer;

#ifdef __cplusplus
extern "C" {
#endif
route_layer make_route_layer(int batch, int n, int *input_layers, int *input_size, int groups, int group_id);
void forward_route_layer(const route_layer l, network_state state);
void backward_route_layer(const route_layer l, network_state state);
void resize_route_layer(route_layer *l, network *net);

#ifdef GPU
void forward_route_layer_gpu(const route_layer l, network_state state);
void backward_route_layer_gpu(const route_layer l, network_state state);
#endif

#ifdef __cplusplus
}
#endif
#endif

其中有幾個比較關鍵的函數聲明，forward_route_layer 和 backward_route_layer 是 cpu 上的前向和後向，forward_route_layer_gpu 和 backward_route_layer_gpu 是 gpu 上的前向和後向。make_route_layer 是用于建構 route 層的，裡面會調 cpu 和 gpu 的前向和後面，當然 gpu 的話需要你開了 gpu 才會去調，resize_route_layer 是經過 route 層作用後資料次元方面的變換。

來看一下 darknet 建構 route 的操作，darkent 裡 src/ 有一個 darknet.c，裡面的 main 是整個架構的入口，裡面也提供了功能豐富的傳參示例能夠幫你快速應用 darkent 這個好用的架構。順着 darknet.c 裡的邏輯你會找到建構 route 層的代碼：

else if(lt == ROUTE){
            l = parse_route(options, params);
            int k;
            for (k = 0; k < l.n; ++k) {
                net.layers[l.input_layers[k]].use_bin_output = 0;
                net.layers[l.input_layers[k]].keep_delta_gpu = 1;
            }

我們來看看 parse_route 做了什麼：

route_layer parse_route(list *options, size_params params)
{
    char *l = option_find(options, "layers");
    if(!l) error("Route Layer must specify input layers");
    int len = strlen(l);
    int n = 1;
    int i;
    for(i = 0; i < len; ++i){
        if (l[i] == ',') ++n;
    }

    int* layers = (int*)xcalloc(n, sizeof(int));
    int* sizes = (int*)xcalloc(n, sizeof(int));
    for(i = 0; i < n; ++i){
        int index = atoi(l);
        l = strchr(l, ',')+1;
        if(index < 0) index = params.index + index;
        layers[i] = index;
        sizes[i] = params.net.layers[index].outputs;
    }
    int batch = params.batch;

    int groups = option_find_int_quiet(options, "groups", 1);
    int group_id = option_find_int_quiet(options, "group_id", 0);

    route_layer layer = make_route_layer(batch, n, layers, sizes, groups, group_id);

    convolutional_layer first = params.net.layers[layers[0]];
    layer.out_w = first.out_w;
    layer.out_h = first.out_h;
    layer.out_c = first.out_c;
    for(i = 1; i < n; ++i){
        int index = layers[i];
        convolutional_layer next = params.net.layers[index];
        if(next.out_w == first.out_w && next.out_h == first.out_h){
            layer.out_c += next.out_c;
        }else{
            fprintf(stderr, " The width and height of the input layers are different. \n");
            layer.out_h = layer.out_w = layer.out_c = 0;
        }
    }
    layer.out_c = layer.out_c / layer.groups;

    layer.w = first.w;
    layer.h = first.h;
    layer.c = layer.out_c;

    if (n > 3) fprintf(stderr, " \t    ");
    else if (n > 1) fprintf(stderr, " \t            ");
    else fprintf(stderr, " \t\t            ");

    fprintf(stderr, "           ");
    if (layer.groups > 1) fprintf(stderr, "%d/%d", layer.group_id, layer.groups);
    else fprintf(stderr, "   ");
    fprintf(stderr, " -> %4d x%4d x%4d \n", layer.out_w, layer.out_h, layer.out_c);

    return layer;
}

其中關鍵的是：

route_layer layer = make_route_layer(batch, n, layers, sizes, groups, group_id);

這就來到了上面 route_layer 頭裡的 make_route_layer 了：

route_layer make_route_layer(int batch, int n, int *input_layers, int *input_sizes, int groups, int group_id)
{
    fprintf(stderr,"route ");
    route_layer l = { (LAYER_TYPE)0 };
    l.type = ROUTE;
    l.batch = batch;
    l.n = n;
    l.input_layers = input_layers;
    l.input_sizes = input_sizes;
    l.groups = groups;
    l.group_id = group_id;
    int i;
    int outputs = 0;
    for(i = 0; i < n; ++i){
        fprintf(stderr," %d", input_layers[i]);
        outputs += input_sizes[i];
    }
    outputs = outputs / groups;
    l.outputs = outputs;
    l.inputs = outputs;
    //fprintf(stderr, " inputs = %d \t outputs = %d, groups = %d, group_id = %d \n", l.inputs, l.outputs, l.groups, l.group_id);
    l.delta = (float*)xcalloc(outputs * batch, sizeof(float));
    l.output = (float*)xcalloc(outputs * batch, sizeof(float));

    l.forward = forward_route_layer;
    l.backward = backward_route_layer;
    #ifdef GPU
    l.forward_gpu = forward_route_layer_gpu;
    l.backward_gpu = backward_route_layer_gpu;

    l.delta_gpu =  cuda_make_array(l.delta, outputs*batch);
    l.output_gpu = cuda_make_array(l.output, outputs*batch);
    #endif
    return l;
}

來看一下這個函數，傳參中 n 是輸入層的個數，groups 和 group_id 就是前面提到的 route 作 slice 的參數，由下面定義可以知道 groups 預設會是 1，group_id 預設會是 0，也就是當 route 沒有 groups 參數時，上面函數中的 outputs = outputs / groups 和 group_id 偏移就相當于啥都沒做。

int groups = option_find_int_quiet(options, "groups", 1);
int group_id = option_find_int_quiet(options, "group_id", 0);

來看一下 forward_route_layer：

void forward_route_layer(const route_layer l, network_state state)
{
    int i, j;
    int offset = 0;
    for(i = 0; i < l.n; ++i){
        int index = l.input_layers[i];
        float *input = state.net.layers[index].output;
        int input_size = l.input_sizes[i];
        int part_input_size = input_size / l.groups;
        for(j = 0; j < l.batch; ++j){
            //copy_cpu(input_size, input + j*input_size, 1, l.output + offset + j*l.outputs, 1);
            copy_cpu(part_input_size, input + j*input_size + part_input_size*l.group_id, 1, l.output + offset + j*l.outputs, 1);
        }
        //offset += input_size;
        offset += part_input_size;
    }
}

最關鍵的實作很簡單就是在 copy_cpu 的一個指派語句，其實想想也是，route 并沒有做什麼複雜的操作，拼接或者切割或者鍊路的操作用一些指針偏移和指派就可以實作了。

void copy_cpu(int N, float *X, int INCX, float *Y, int INCY)
{
    int i;
    for(i = 0; i < N; ++i) Y[i*INCY] = X[i*INCX];
}

同樣，高性能計算離不開 cuda，來看一下 forward_route_layer_gpu 的操作：

void forward_route_layer_gpu(const route_layer l, network_state state)
{
    if (l.stream >= 0) {
        switch_stream(l.stream);
    }

    if (l.wait_stream_id >= 0) {
        wait_stream(l.wait_stream_id);
    }

    int i, j;
    int offset = 0;
    for(i = 0; i < l.n; ++i){
        int index = l.input_layers[i];
        float *input = state.net.layers[index].output_gpu;
        int input_size = l.input_sizes[i];
        int part_input_size = input_size / l.groups;
        for(j = 0; j < l.batch; ++j){
            //copy_ongpu(input_size, input + j*input_size, 1, l.output_gpu + offset + j*l.outputs, 1);
            //simple_copy_ongpu(input_size, input + j*input_size, l.output_gpu + offset + j*l.outputs);
            simple_copy_ongpu(part_input_size, input + j*input_size + part_input_size*l.group_id, l.output_gpu + offset + j*l.outputs);
        }
        //offset += input_size;
        offset += part_input_size;
    }
}

route_gpu 的操作邏輯和 route_cpu 的一樣，差別在于 gpu 的資料拷貝給并行起來了，來看一下 simple_copy_ongpu：

extern "C" void simple_copy_ongpu(int size, float *src, float *dst)
{
    const int num_blocks = size / BLOCK + 1;
    simple_copy_kernel << <num_blocks, BLOCK, 0, get_cuda_stream() >> >(size, src, dst);
    CHECK_CUDA(cudaPeekAtLastError());
}

simple_copy_ongpu 裡調了 simple_copy_kernel 這個 cuda kernel，我們來看一下：

__global__ void simple_copy_kernel(int size, float *src, float *dst)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    if (index < size)
        dst[index] = src[index];
}

很簡單，就是一個 device_to_device 的資料指派過程。

說到這裡應該是比較清晰的了，不得不感慨一句：C 語言看起來真清爽啊~

好了，有問題歡迎讨論，收工~

【公衆号傳送】

【模型推理】談談 darknet yolo 的 route 算子

文章目錄

1、route 化作 concatenate

2、route 化作 slice

3、route 化作恒等映射

4、darknet 源碼解讀

繼續閱讀

241 Different Ways to Add Parentheses（C代碼版）

【趨高機器視覺】機器視覺技術原了解析及解決方案

CSMA/CD1． CSMA/CD的概述2． CSMA 的工作原理3． CSMA/CD控制規程及特點4． CSMA/CD協定5． CSMA/CD的優點6．結束語

極大似然法(ML)與最大期望法(EM)

[HTML5]自定義屬性 data-* 和 jQuery.data 詳解

解碼器用于語義分割：資料依賴的解碼可以實作靈活的特征聚合

2021-2025年中國運動療法（KT）帶行業市場供需與戰略研究報告

cs231n斯坦福基于卷積神經網絡的CV學習筆記（一）KNN和線性分類器/分類器損失/反向傳播一，KNN圖像分類算法二，線性分類器三，線性分類器損失四，反向傳播五，神經網絡

C++ 第十五周報告1--《冒泡法排序》

2021年危險化學品經營機關安全管理人員考試題庫及危險化學品經營機關安全管理人員考試技巧

筆試面試題目：滑動視窗(二)

資料結構與算法（27）——排序（二）

無人機--飛控科普

Dijkstra--簡易版（最短路徑）

GitHub連夜封殺！這份阿裡 10W 字内部 Java 字面試手冊到底有多強？

hdu7108哈希