CAFFE源码学习笔记之四-device_alternate

一、前言
common中包含的头文件中有一个device_alternate,里面主要是和cuda有关系的宏：
CUDA_CHECK;
CUBLAS_CHECK;
CURAND_CHECK;
CUDA_KERNEL_LOOP;
同时还对block中的线程数以及block的维度声明为常量。
二、源码分析
首先是NO_GPU宏，主要是当设置为CPU_ONLY后，如果你仍然固执地要用forward的gpu版，那么他就会告诉你不要犯傻好么。LOG(FATAL)是glog的宏，后面集中总结一下。

#define NO_GPU LOG(FATAL) << "Cannot use GPU in CPU-only Caffe: check mode."

cudaError_t是cuda 运行时API，任何cuda API运行都有可能发生错误，只有当正常运行后才返回cudaSuccess。为什么专门写宏呢，为了避免cudaError_t重复定义，同时也是省事。

#define CUDA_CHECK(condition) \
  /* Code block avoids redefinition of cudaError_t error */ \
  do { \
    cudaError_t error = condition; \
    CHECK_EQ(error, cudaSuccess) << " " << cudaGetErrorString(error); \
  } while (0)

const定义block内的线程个数为512，而block的个数表示为数据点总数除以每个block内的线程数，就是每个线程负责一个计算。

// CUDA: use 512 threads per block
const int CAFFE_CUDA_NUM_THREADS = 512;

// CUDA: number of blocks for threads.
inline int CAFFE_GET_BLOCKS(const int N) {
  return (N + CAFFE_CUDA_NUM_THREADS - 1) / CAFFE_CUDA_NUM_THREADS;

循环以线程总数为步长，每次循环保证所有线程同时进行计算。

// 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)

四、总结
主要是写宏可以很省事。至于cuda再开一系列。。。