caffe源码分析--math_functions.cu代码研究

其中用到一个宏定义CUDA_KERNEL_LOOP
在common.hpp中有。

#defineCUDA_KERNEL_LOOP(i,n)
\
for(inti
= blockIdx.x * blockDim.x + threadIdx.x; \
i < (n); \
i +=blockDim.x * gridDim.x)

先看看caffe采取的线程格和线程块的维数设计，
还是从common.hpp可以看到
CAFFE_CUDA_NUM_THREADS
CAFFE_GET_BLOCKS(constintN)
明显都是一维的。

整理一下CUDA_KERNEL_LOOP格式看看，
for(inti
= blockIdx.x * blockDim.x + threadIdx.x;
i< (n);
i+= blockDim.x * gridDim.x)
blockDim.x* gridDim.x表示的是该线程格所有线程的数量。
n表示核函数总共要处理的元素个数。
有时候，n会大于blockDim.x*
gridDim.x，因此并不能一个线程处理一个元素。
由此通过上面的方法，让一个线程串行（for循环）处理几个元素。
这其实是常用的伎俩，得借鉴学习一下。

再来看一下这个核函数的实现。

template<typename Dtype>
__global__void mul_kernel(const int n, const Dtype* a,
constDtype* b, Dtype* y)
{
CUDA_KERNEL_LOOP(index,n)
{
y[index]= a[index] * b[index];
}
}

明显就是算两个向量的点积了。
由于向量的维数可能大于该kernel函数线程格的总线程数量。
因此有些线程可以要串行处理几个元素。