基于Caffe的Large Margin Softmax Loss的实现(中)

小喵的唠叨话:前一篇博客,我们做完了L-Softmax的准备工作。而这一章,我们开始进行前馈的研究。 四、前馈 还记得上一篇博客,小喵给

小喵的唠叨话:前一篇博客,我们做完了L-Softmax的准备工作。而这一章,我们开始进行前馈的研究。

四、前馈

还记得上一篇博客,小喵给出的三个公式吗?不记得也没关系。

这次,我们要一点一点的通过代码来实现这些公式。小喵主要是GPU上实现前后馈的代码,因为这个层只是用来训练,GPU速度应该会快一点。

我们首先要进行一般的FC层的前馈,因为LM_FC的前馈只是修改了一般的FC中的若干个值,而大部分的值都是没有修改过的。

const Dtype* bottom_data = bottom[0]->gpu_data();const Dtype* label_data = bottom[1]->gpu_data();Dtype* top_data = top[0]->mutable_gpu_data();const Dtype* weight = this->blobs_[0]->gpu_data();// 普通fc层的计算if (M_ == 1) {  caffe_gpu_gemv<Dtype>(CblasNoTrans, N_, K_, (Dtype)1.,                      weight, bottom_data, (Dtype)0., top_data);} else {  caffe_gpu_gemm<Dtype>(CblasNoTrans,                        transpose_ ? CblasNoTrans : CblasTrans,                        M_, N_, K_, (Dtype)1.,                        bottom_data, weight, (Dtype)0., top_data);}

这样就计算完了一个普通的FC的前馈。

之后是一些具体的实现。

1, 基于Caffe的Large Margin Softmax Loss的实现(中)

这是要求出label为的weight的权值和feature之间的余弦值。公式大家在高中应该就学过了。这样需要出三部分: 基于Caffe的Large Margin Softmax Loss的实现(中)基于Caffe的Large Margin Softmax Loss的实现(中) 和。这里表示feature的序号,因为一个mini batch中有很多张图片。表示正确的label值。

基于Caffe的Large Margin Softmax Loss的实现(中) 的计算非常简单,因为FC层的前馈计算出来的就是这个值。因此我们可以直接从FC的前馈结果中直接复制对应位置的结果。

基于Caffe的Large Margin Softmax Loss的实现(中) 和是比较简单的模值的计算,使用caffe_cpu_dot很容易就可以求得(为什么不使用caffe_gpu_dot呢?因为小喵在使用caffe_gpu_dot的时候,caffe会报一个奇怪的错误,不知道是不是因为GPU的显存不能随意访问的)。

最后的余弦值带入到上面的式子,就一下子搞定~

这里用到了几个变量:

M_: batch size

N_: class num

K_: feature length

// w * x// 直接从前馈的结果中复制Dtype *wx_data = this->wx_.mutable_gpu_data();copy_label_score<Dtype><<<CAFFE_GET_BLOCKS(M_), CAFFE_CUDA_NUM_THREADS>>>(M_, N_, label_data, top_data, wx_data); // w * wDtype *abs_w_data = this->abs_w_.mutable_cpu_data();for (int m = 0; m < M_; ++ m) {  abs_w_data[m] = caffe_cpu_dot<Dtype>(    K_,    this->blobs_[0]->cpu_data() + static_cast<int>(label_cpu_data[m]) * K_,    this->blobs_[0]->cpu_data() + static_cast<int>(label_cpu_data[m]) * K_    );} // x * xDtype *abs_x_data = this->abs_x_.mutable_cpu_data();for (int m = 0; m < M_; ++ m) {  abs_x_data[m] = caffe_cpu_dot<Dtype>(    K_,     bottom[0]->cpu_data() + m * K_,    bottom[0]->cpu_data() + m * K_    );}// abs_w, abs_xcaffe_gpu_powx<Dtype>(M_, this->abs_w_.mutable_gpu_data(), 0.5, this->abs_w_.mutable_gpu_data());caffe_gpu_powx<Dtype>(M_, this->abs_x_.mutable_gpu_data(), 0.5, this->abs_x_.mutable_gpu_data()); // cos_t = wx / (|x| * |w|)Dtype *cos_t_data = this->cos_t_.mutable_gpu_data();caffe_gpu_div<Dtype>(M_, wx_data, this->abs_x_.gpu_data(), cos_t_data);caffe_gpu_div<Dtype>(M_, cos_t_data, this->abs_w_.gpu_data(), cos_t_data);

其中copy_label_score是我们自己编写的用来复制结果的核函数(如何编写Cuda程序就是另一门学科了):

template <typename Dtype>__global__void copy_label_score(const int M, const int N, const Dtype *label_data, const Dtype *top_data, Dtype *wx_data) {  CUDA_KERNEL_LOOP(index, M) {    wx_data[index] = top_data[index * N + static_cast<int>(label_data[index])];  }}

相信机智如你的喵粉,看到这几行代码,一定可以轻松理解。

这里,小喵想多介绍一点东西。

我们知道Caffe里面的数据都是通过Blob结构来存储的,比如这里的bottom_data,其实就是一个blob,默认形状是(n, c, h, w),n表示的就是batch size,c是channel数,h,w分贝表示高和宽。而且blob中的内存的存储顺序,也和一般的C语言中的数组一样。因此我们这里计算feature的模的时候,是直接每K_个数值计算一次点乘。

同理,weight是存储在this->blobs[0]中的,那么weight的形状又是什么样子的呢?这里非常碰巧的是,如果我们在prototxt中设置的transpose为false的话,weight的形状是N_*K_,也就是说,我们可以将weight看成一个矩阵,它的每一行都与feature直接点乘,得到输出,也就是说weight的每一行都是我们需要计算模值的,所以我们计算weight的模的时候,用的计算方法和计算feature模时很相似。我们这里强制设置transpose为false,因为这样计算会比较简单。如果你设成了true,那就必须自己写个求模的函数了。

2, 基于Caffe的Large Margin Softmax Loss的实现(中)

我们在(1)中求出了,对于给定的margin,只需要代入公式就可以求出的值了。

template <typename Dtype>__global__void cal_cos_mt(const int count, const unsigned int margin, const int *C_M_N, const Dtype *cos_t_data, Dtype *cos_mt_data) {  CUDA_KERNEL_LOOP(index, count) {    Dtypecos_t = cos_t_data[index];    Dtypesin_t_2 = 1 - cos_t * cos_t;    Dtypecos_mt = 0.;    int flag = -1;    for (int n = 0; n <= (margin / 2); ++ n) {      flag *= -1;      cos_mt += flag * C_M_N[2 * n] * powf(cos_t, (margin - 2 * n)) * powf(sin_t_2, n);    }    cos_mt_data[index] = cos_mt;  }}

上面是用来计算的cuda函数,调用也十分的简单:

// cos(mt)cal_cos_mt<Dtype><<<CAFFE_GET_BLOCKS(M_), CAFFE_CUDA_NUM_THREADS>>>(  M_, this->margin, this->C_M_N_.gpu_data(), this->cos_t_.mutable_gpu_data(), this->cos_mt_->mutable_gpu_data());

3, 基于Caffe的Large Margin Softmax Loss的实现(中)

严格上来说,我们需要求的并不是这个式子,而是:

  基于Caffe的Large Margin Softmax Loss的实现(中)

  基于Caffe的Large Margin Softmax Loss的实现(中)

可以看出,当为0的时候,这两个式子就退化成前面的一个式子了。

k的求法十分简单,只需要将与各个区间进行比较就可以得到。

// kint *k_cpu_data = this->k_.mutable_cpu_data();const Dtype *cos_t_cpu_data = this->cos_t_.cpu_data();for (int m = 0; m < M_; ++ m) {  for (int _k = 0; _k < this->cos_theta_bound_.count(); ++ _k) {    if (this->cos_theta_bound_.cpu_data()[_k] < cos_t_cpu_data[m]) {      k_cpu_data[m] = _k - 1;      break;    }  }}

最后一步就是计算出真正的前馈值了!按照公式容易编写程序:

template <typename Dtype>__global__void LMForward(  const int M, const int N, const float lambda,  const Dtype *label_data, const Dtype *cos_mt_data, const int *k_data,  const Dtype *abs_w_data, const Dtype *abs_x_data, Dtype *top_data) {   CUDA_KERNEL_LOOP(index, M) {    Dtypecos_mt = cos_mt_data[index];    int k = k_data[index];    int label = static_cast<int>(label_data[index]);    Dtypeabs_w = abs_w_data[index];    Dtypeabs_x = abs_x_data[index];    top_data[N * index + label] =  (lambda * top_data[N * index + label] + abs_w * abs_x * ( powf(-1, k) * cos_mt - 2 * k )) / (1 + lambda);  }}

调用也十分简单:

// yLMForward<Dtype><<<CAFFE_GET_BLOCKS(M_), CAFFE_CUDA_NUM_THREADS>>>(  M_, N_, this->lambda,  label_data, this->cos_mt_->gpu_data(), this->k_.gpu_data(),  this->abs_w_.gpu_data(), this->abs_x_.gpu_data(), top[0]->mutable_gpu_data());

最后附上,完整的前馈代码(省略头文件和caffe的名字空间):

template <typenameDtype>__global__void copy_label_score(const int M, const int N, const Dtype *label_data, const Dtype *top_data, Dtype *wx_data) {  CUDA_KERNEL_LOOP(index, M) {    wx_data[index] = top_data[index * N + static_cast<int>(label_data[index])];  }} template <typenameDtype>__global__void cal_cos_mt(const int count, const unsigned int margin, const int *C_M_N, const Dtype *cos_t_data, Dtype *cos_mt_data) {  CUDA_KERNEL_LOOP(index, count) {    Dtypecos_t = cos_t_data[index];    Dtypesin_t_2 = 1 - cos_t * cos_t;    Dtypecos_mt = 0.;    int flag = -1;    for (int n = 0; n <= (margin / 2); ++ n) {      flag *= -1;      cos_mt += flag * C_M_N[2 * n] * powf(cos_t, (margin - 2 * n)) * powf(sin_t_2, n);    }    cos_mt_data[index] = cos_mt;  }} template <typenameDtype>__global__void LMForward(  const int M, const int N, const float lambda,  const Dtype *label_data, const Dtype *cos_mt_data, const int *k_data,  const Dtype *abs_w_data, const Dtype *abs_x_data, Dtype *top_data) {   CUDA_KERNEL_LOOP(index, M) {    Dtypecos_mt = cos_mt_data[index];    int k = k_data[index];    int label = static_cast<int>(label_data[index]);    Dtypeabs_w = abs_w_data[index];    Dtypeabs_x = abs_x_data[index];    top_data[N * index + label] =  (lambda * top_data[N * index + label] + abs_w * abs_x * ( powf(-1, k) * cos_mt - 2 * k )) / (1 + lambda);  }} template <typenameDtype>void LargeMarginInnerProductLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,    const vector<Blob<Dtype>*>& top) {  const Dtype* bottom_data = bottom[0]->gpu_data();  const Dtype* label_data = bottom[1]->gpu_data();  Dtype* top_data = top[0]->mutable_gpu_data();  const Dtype* weight = this->blobs_[0]->gpu_data();   // 普通fc层的计算  if (M_ == 1) {    caffe_gpu_gemv<Dtype>(CblasNoTrans, N_, K_, (Dtype)1.,                        weight, bottom_data, (Dtype)0., top_data);  } else {    caffe_gpu_gemm<Dtype>(CblasNoTrans,                          transpose_ ? CblasNoTrans : CblasTrans,                          M_, N_, K_, (Dtype)1.,                          bottom_data, weight, (Dtype)0., top_data);  }   const Dtype* label_cpu_data = bottom[1]->cpu_data();   // w * x  // 直接从前馈的结果中复制  Dtype *wx_data = this->wx_.mutable_gpu_data();  copy_label_score<Dtype><<<CAFFE_GET_BLOCKS(M_), CAFFE_CUDA_NUM_THREADS>>>(M_, N_, label_data, top_data, wx_data);   // w * w  Dtype *abs_w_data = this->abs_w_.mutable_cpu_data();  for (int m = 0; m < M_; ++ m) {    abs_w_data[m] = caffe_cpu_dot<Dtype>(      K_,      this->blobs_[0]->cpu_data() + static_cast<int>(label_cpu_data[m]) * K_,      this->blobs_[0]->cpu_data() + static_cast<int>(label_cpu_data[m]) * K_      );  }    // x * x  Dtype *abs_x_data = this->abs_x_.mutable_cpu_data();  for (int m = 0; m < M_; ++ m) {    abs_x_data[m] = caffe_cpu_dot<Dtype>(      K_,       bottom[0]->cpu_data() + m * K_,      bottom[0]->cpu_data() + m * K_      );  }   // abs_w, abs_x  caffe_gpu_powx<Dtype>(M_, this->abs_w_.mutable_gpu_data(), 0.5, this->abs_w_.mutable_gpu_data());  caffe_gpu_powx<Dtype>(M_, this->abs_x_.mutable_gpu_data(), 0.5, this->abs_x_.mutable_gpu_data());   // cos_t = wx / (|x| * |w|)  Dtype *cos_t_data = this->cos_t_.mutable_gpu_data();  caffe_gpu_div<Dtype>(M_, wx_data, this->abs_x_.gpu_data(), cos_t_data);  caffe_gpu_div<Dtype>(M_, cos_t_data, this->abs_w_.gpu_data(), cos_t_data);   // cos(mt)  cal_cos_mt<Dtype><<<CAFFE_GET_BLOCKS(M_), CAFFE_CUDA_NUM_THREADS>>>(    M_, this->margin,     this->C_M_N_.gpu_data(),     this->cos_t_.gpu_data(),    this->cos_mt_.mutable_gpu_data()    );   // k  int *k_cpu_data = this->k_.mutable_cpu_data();  const Dtype *cos_t_cpu_data = this->cos_t_.cpu_data();  for (int m = 0; m < M_; ++ m) {    for (int _k = 0; _k < this->cos_theta_bound_.count(); ++ _k) {      if (this->cos_theta_bound_.cpu_data()[_k] < cos_t_cpu_data[m]) {        k_cpu_data[m] = _k - 1;        break;      }    }  }   // y  LMForward<Dtype><<<CAFFE_GET_BLOCKS(M_), CAFFE_CUDA_NUM_THREADS>>>(    M_, N_, this->lambda,    label_data, this->cos_mt_.gpu_data(), this->k_.gpu_data(),    this->abs_w_.gpu_data(), this->abs_x_.gpu_data(), top[0]->mutable_gpu_data());}

那么,这样关于large margin softmax loss的前馈我们就轻松的实现了。下一篇,我们要讲最复杂的后馈的实现了。

如果您觉得本文对您有帮助,那请小喵喝杯茶吧~~O(∩_∩)O~~ 再次感慨大法好。

基于Caffe的Large Margin Softmax Loss的实现(中)

转载请注明出处~

未登录用户
全部评论0
到底啦