1. Cuda中要处理单位数据N大于可用的线程数量N’时
以向量乘函数为例,mul_kernel(n,a,b,y)对长为n的a,b求内积,结果放入y
template <typename Dtype>
__global__ void mul_kernel(const int n, const Dtype* a,
const Dtype* b, Dtype* y) {
CUDA_KERNEL_LOOP(index, n) {
y[index] = a[index] * b[index];
}
}
展开CUDA_KERNEL_LOOP(index, n)得
template <typename Dtype>
__global__ void mul_kernel(const int n, const Dtype* a,
for (int i = blockIdx.x * blockDim.x + threadIdx.x; \
i < (n); \
i += blockDim.x * gridDim.x)
y[i] = a[i] * b[i];
}
}
说明:
- Caffe中,Grid、Block都是一维的,其中Block的维度为512(也就是一个Block有512个Thread)
- n为单位总数
- blockDim.x* gridDim.x为该Grid所有线程的数量
- 可见,一个线程可能处理了多个单位