caffe下relu_layer.cu详解

最新推荐文章于 2019-01-10 15:21:54 发布

原创最新推荐文章于 2019-01-10 15:21:54 发布 · 1.4k 阅读

0 ·

CC 4.0 BY-SA版权

文章标签：

#gpu #caffe源码

cuda编程同时被 2 个专栏收录

2 篇文章

订阅专栏

caffe

1 篇文章

订阅专栏

本文深入解析了Caffe框架中ReLU层的GPU实现原理。重点介绍了ReLULayer::Forward_gpu函数的工作流程，并解释了其中ReLUForward内核的具体实现细节及宏定义的作用，帮助读者理解为何需要循环处理数据。

relu_layer是caffe框架的一个线性激活单元，具体功能、作用、和c++代码我不做详细分析，相信有点c++基础和深度学习基础的孩子都能看懂，今天我来详细分析relu_layer.cu文件，因为里面有点小坑思考了很久才想明白，写出来分享一波。
首先是ReLULayer<Dtype>::Forward_gpu函数，这是c++写的代码，具体难度应该不大。

template <typename Dtype>
void ReLULayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
    const vector<Blob<Dtype>*>& top) {
  const Dtype* bottom_data = bottom[0]->gpu_data();
  Dtype* top_data = top[0]->mutable_gpu_data();
  const int count = bottom[0]->count();
  Dtype negative_slope = this->layer_param_.relu_param().negative_slope();
  // NOLINT_NEXT_LINE(whitespace/operators)
  ReLUForward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
      count, bottom_data, top_data, negative_slope);
  CUDA_POST_KERNEL_CHECK;
  // << " count: " << count << " bottom_data: "
  //     << (unsigned long)bottom_data
  //     << " top_data: " << (unsigned long)top_data
  //     << " blocks: " << CAFFE_GET_BLOCKS(count)
  //     << " threads: " << CAFFE_CUDA_NUM_THREADS;
}

ReLUForward<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>(
      count, bottom_data, top_data, negative_slope);

追踪这个函数可以看到代码是这样

template <typename Dtype>
__global__ void ReLUForward(const int n, const Dtype* in, Dtype* out,
    Dtype negative_slope) {
  CUDA_KERNEL_LOOP(index, n) {
    out[index] = in[index] > 0 ? in[index] : in[index] * negative_slope;
  }
}

里面出现了宏定义，可以跳转宏定义查看源代码

// CUDA: grid stride looping
#define CUDA_KERNEL_LOOP(i, n) \
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
       i < (n); \
       i += blockDim.x * gridDim.x)

这里有个问题，为什么需要这个循环呢？不是一个线程处理一个数据吗？
这个循环其实是为了防止用户没有分配足够线程数而准备的，使每个线程串行处理多个数据，当用户分配足够多的线程时其实循环是不起任何作用的。例如 caffe上面的代码，分配的线程是足够的，所以循环是不起任何作用的。这样写应该是为了确保安全可靠性。
如果分析不正确或者有误，还望赐教。