caffe中能够很方便地切换CPU或者GPU模式，但是如果我们想要develop 新的层，则需要写cpp和cu文件，相比之下theano比caffe的扩展方便地多。首次读Caffe的代码，简单记录学习下cu文件如何实现。

CUDA C

CUDA C最简单的形式就是C语言而已，如果不涉及GPU的操作，当然这样也还是能够被NVIDIA的编译器NVCC所编译。

__global__ void kernel(void){}

CUDA 中 __global__ 表明该函数在device(GPU)设备上运行，并且由host(CPU)调用。例如一个加法程序

__global__ void add(int* a, int* b, int*c)
{
    *c = *a + *b;
}
int main(void)
{
    int a,b,c;
    int *gpu_ptr_a,gpu_ptr_b,gpu_ptr_c;
    int size = sizeof(int);
    cudaMalloc((void**)&gpu_ptr_a,size);
    cudaMalloc((void**)&gpu_ptr_b,size);
    cudaMalloc((void**)&gpu_ptr_c,size);
    a = 2;
    b = 7;
    cudaMemcpy(gpu_ptr_a,&a,size,cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr_b,&b,size,cudaMemcpyHostToDevice);

    add<<<1,1>>>(gpu_ptr_a,gpu_ptr_b,gpu_ptr_c);

    cudaMemcpy(&c,gpu_ptr_c,size,cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr_a);
    cudaFree(gpu_ptr_b);
    cudaFree(gpu_ptr_c);
    return 0;
}

内存管理

CUDA对device(GPU )的内存管理主要通过cudaMalloc()、cudaFree()、cudaMemcpy() 进行管理。另外，从上述代码我们可以看到，add() 函数的调用比较奇怪相对于C语言来说，需要用add<<<M，N>>> 这种形式表明这是一个从host(CPU)代码调用device的代码，并且括号中的数值表明，M个block，每个block有 N个线程, 所以这个函数总共有M*N个线程。

CUDA 并行操作

CUDA中用来实现并行操作的有block 和thread 两个模块。

• block在代码中用blockIdx.x 指示。blockIdx.x为cuda中内建(build-in)的变量，它表明正在执行状态下的block的index。Cuda允许使用多维的索引，.x 为常用的一维索引。一个block可以切分成不同的threads。

• thread 在代码中用 threadIdx.x表示线程的标号。thread是block的进一步划分，thread好处在于并行的block难以进行通信和同步(Communciate and Synchronize )，thread通过Shared Memory 在线程之间进行同步，在CUDA C用关键字__shared__ 表示。

  #define N 512
  __global__ void dot(int *a,int *b, int *c)
  {
    __shared__ int temp[N];
    temp[threadIdx.x] = a[threadIdx.x] * b[threadIdx.x];
  
    __syncthreads();
    if(0 == threadIdx.x)
    {
        int sum = 0;
        for(int i=0; i<N; i++)
        {
            sum += temp[i] ;
        }
        *c = sum;
    }
  }

  上述代码实现了向量点积的过程。 涉及到多线程操作和内存共享的问题，必须要考虑线程之间的同步。__syncthreads() 实现的就是这个目的，只有当所有的线程运行到了__syncthreads() 处(即共享的temp[N]已经运算完成)线程才能继续向下执行，否者temp[N] 还没写完就去读，sum得到的结果就未知了。

  sigmoid_layer.cu 简析

  cu文件的组成:

• host(CPU)调用的函数Forward_gpu()和Backward_gpu()

• 在device(GPU)上计算的核函数(由CPU中的函数调用)

• 实例化函数模板,包括Forward_gpu()和Backward_gpu()

下面以sigmoid_layer的CUDA实现进行简要分析。

template <typename Dtype>
__global__ void SigmoidForward(const int n, const Dtype* in, Dtype* out) {
  CUDA_KERNEL_LOOP(index, n) {
    out[index] = 1. / (1. + exp(-in[index]));
  }
}

上面是sigmoid的前向传播函数的核函数，这里CUDA_KERNEL_LOOP其实是定义在在device_alternate.hpp 的宏，CUDA_KERNEL_LOOP的详细解释可以参考这篇问答

#define CUDA_KERNEL_LOOP(i, n) \
  for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
       i < (n); \
       i += blockDim.x * gridDim.x)

但是caffe中与此处 的解释不太相同。caffe中固定每个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;
}

所以利用CAFFE_GET_BLOCKS 方法，线程的个数始终大于或等于元素个数N，故不存在一个线程 for循环处理多个元素。在CUDA_KERNEL_LOOP 中，只有在线程的个数小于元素个数N的情况下for循环才起效。

Forward_gpu() 、 Backward_gpu() 调用核函数实现就不展开介绍了。

最后cu文件还需要对Forward_gpu()和Backward_gpu() 进行显示实例化(instantiation)。

INSTANTIATE_LAYER_GPU_FUNCS(SigmoidLayer);

上述代码是Caffe中定义的一个宏，具体展开如下

#define INSTANTIATE_LAYER_GPU_FORWARD(classname) \
template void classname<float>::Forward_gpu( \
const std::vector<Blob<float>*>& bottom, \
const std::vector<Blob<float>*>& top); \
template void classname<double>::Forward_gpu( \
const std::vector<Blob<double>*>& bottom, \
const std::vector<Blob<double>*>& top);

#define INSTANTIATE_LAYER_GPU_BACKWARD(classname) \
template void classname<float>::Backward_gpu( \
    const std::vector<Blob<float>*>& top, \
    const std::vector<bool>& propagate_down, \
    const std::vector<Blob<float>*>& bottom); \
template void classname<double>::Backward_gpu( \
    const std::vector<Blob<double>*>& top, \
    const std::vector<bool>& propagate_down, \
    const std::vector<Blob<double>*>& bottom)

 #define INSTANTIATE_LAYER_GPU_FUNCS(classname) \
 INSTANTIATE_LAYER_GPU_FORWARD(classname); \
 INSTANTIATE_LAYER_GPU_BACKWARD(classname)

caffe中定义的都是函数模板，是不会参与编译的，所以只有把函数实例化(INSTANTIATE)，编译器才会编译函数模板。类似的，我们在sigmoid_layer.cpp 文件中，代码最后的 INSTANTIATE_CLASS(SigmoidLayer); 也是同样的道理，将模板类实例化。

//------common.hpp
#define INSTANTIATE_CLASS(classname) \
  char gInstantiationGuard##classname; \
  template class classname<float>; \
  template class classname<double>

reference

• [pdf] GPU Technology conference----CUDA C
• [pdf] CUDA范例精解通用GPU编程
• caffe CUDA 相关宏定义
• under-the-hood-caffe
• sigmoid_layer.cu
• caffe 中的一些编程规范