当中用到一个宏定义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函数线程格的总线程数量。

因此有些线程能够要串行处理几个元素。