当中用到一个宏定义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函数线程格的总线程数量。
因此有些线程能够要串行处理几个元素。