smooth_L1_loss_layer.cu解讀 caffe源碼初認識

.cpp是cpu上運行的代碼，.cu是gpu上運行的代碼。

這是smooth_L1_loss_layer.cu的前向傳播部分

#include "caffe/fast_rcnn_layers.hpp"

namespace caffe {

template <typename Dtype>
__global__ void SmoothL1Forward(const int n, const Dtype* in, Dtype* out) {
  // f(x) = 0.5 * x^2    if |x| < 1
  //        |x| - 0.5    otherwise
  CUDA_KERNEL_LOOP(index, n) {
    Dtype val = in[index];
    Dtype abs_val = abs(val);
    if (abs_val < 1) {
      out[index] = 0.5 * val * val;
    } else {
      out[index] = abs_val - 0.5;
    }
  }
}

template <typename Dtype>
void SmoothL1LossLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
    const vector<Blob<Dtype>*>& top) {
  int count = bottom[0]->count();
  caffe_gpu_sub(
      count,
      bottom[0]->gpu_data(),
      bottom[1]->gpu_data(),
      diff_.mutable_gpu_data());    // d := b0 - b1
  if (has_weights_) {
    caffe_gpu_mul(
        count,
        bottom[2]->gpu_data(),
        diff_.gpu_data(),
        diff_.mutable_gpu_data());  // d := w * (b0 - b1)
  }
  SmoothL1Forward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
      count, diff_.gpu_data(), errors_.mutable_gpu_data());
  CUDA_POST_KERNEL_CHECK;

  Dtype loss;
  caffe_gpu_asum(count, errors_.gpu_data(), &loss);
  top[0]->mutable_cpu_data()[0] = loss / bottom[0]->num();　　　　　　　　　　　　　　注意：這里是bottom[0]->num()，不是bottom[0]->count()
}

blob的主要變量：

shared_ptr<SyncedMemory> data_;
shared_ptr<SyncedMemory> diff_;
vector<int> shape_;
int count_;
int capacity_;

 

blob只是一個基本的數據結構，因此內部的變量相對較少，首先是data_指針，指針類型是shared_ptr，屬於boost庫的一個智能指針，這一部分主要用來申請內存存儲data，data主要是正向傳播的時候用的。同理，diff_主要用來存儲偏差，shape_都是存儲Blob的形狀，count表示Blob中的元素個數，也就是個數*通道數*高度*寬度,capacity表示當前的元素個數，因為Blob可能會reshape。count是一個迭代期參與的圖片個數。帶data的里面存儲的是激活值和W、b，diff中存儲的是殘差和dW、db。

 

blob中除了數據成員之外，也有很多用於操作數據的函數成員，下面就說幾個比較重要的：

    void Blob<Dtype>::Reshape()：這個函數是在原來分配的內存不夠的情況下重新分配內存。

    const Dtype* Blob<Dtype>::cpu_data()：這個是獲取Blob結構體中的data_數據的指針，同時限制不能對返回的指針指向的內容進行更改。

    const Dtype* Blob<Dtype>::cpu_diff()：這個是獲取Blob結構體中的diff_數據的指針，同時限制不能對返回的指針指向的內容進行更改。

    Dtype* Blob<Dtype>::mutable_cpu_data()：獲取Blob結構體中的data_數據的指針，同時可以對指針指向的內容更改。

    Dtype* Blob<Dtype>::mutable_cpu_diff()：獲取Blob結構體中的diff_數據的指針，同時可以對指針指向的內容更改。

    void Blob<Dtype>::ShareData(const Blob& other)：讓其他Blob的data_數據和當前Blob共享。

    void Blob<Dtype>::ShareDiff(const Blob& other)：讓其他Blob的diff_和當前的Blob共享。   

 

blob類里面有重載很多個count()函數，主要還是為了統計blob的容量（volume），或者是某一片（slice），從某個axis到具體某個axis的shape乘積。

inline int count(int start_axis, int end_axis)

int count = bottom[0]->count();      count()沒帶參數，計算的是bottom[0]這個輸入blob所有的元素個數。這里就是計算一個迭代期的所有圖片的所有通道的所有坐標點形成的blob數據結構元素的個數。

top[0]->mutable_cpu_data()[0] = loss / bottom[0]->num();　　num()是計算一個迭代期參與的所有圖片的個數。這里就是求一個迭代期所有幾張圖片的平均loss。            

 

caffe_gpu_asum(count, errors_.gpu_data(), &loss);　　caffe_gpu_asum是對向量進行L1范數計算，實際上就是對向量求其每個元素絕對值的和。第一個參數是要計算的元素的個數。

 

 caffe_gpu_sub(
      count,
      bottom[0]->gpu_data(),
      bottom[1]->gpu_data(),
      diff_.mutable_gpu_data());    // d := b0 - b1
  if (has_weights_) {
    caffe_gpu_mul(
        count,
        bottom[2]->gpu_data(),
        diff_.gpu_data(),
        diff_.mutable_gpu_data());  // d := w * (b0 - b1)
}

caffe_gpu_sub,caffe_gpu_mul:這兩個函數分別實現element-wise（即點乘，每個矩陣對應元素相乘）的乘減（y[i] = a[i] * - b[i]）。第一個參數是要計算的元素個數。

 

總結;   smooth_L1_loss_layer的loss計算是將所有對應元素（某張圖片，某個通道的對應坐標）相減，判斷絕對值是否小於1然后各個元素分別進行smooth_L1(x)這個函數的處理，各個元素都有一個loss，然后把所有的loss相加除以圖片數，就得到每張圖片box_loss的值。      

 

loss的兩個輸入是1x84維的向量（fast中是這樣，faster中的rpn是36*w*h），這個向量表示21類的dx，dy，dh，dw。count數出所有的個數，然后兩個輸入相對應的每一個進行這個計算，計算出84個loss，再對84個loss求和。當然這是單個圖片，如果batch有多個圖片，對多個圖片loss求平均。

fast中使用的smoothL1和faster中使用的smoothL1有一點差別，但不大。faster中除了在rpn使用smoothl1，還要在fast那部分使用，所以faster中的smoothl1應該是兼容的。