由于GPU硬件的限制,核函数kernel<<<B,T>>>中B的数量限制为不超过65535;同样对于启动核函数中每个线程块中的线程数
量,T不能超过设备属性结构(cuDeviceProp)中maxThreadsPerBlock的值,对于目前的GPU,该限制值是每个线程块512个线程。即
当矢量的长度超过65536*128时,核函数的调用会失败,而目前的GPU对处理这种量级的运算时很常见的。
当矢量求和计算量大于B*T时,可以在每个线程中加一个B*T的偏移,相当于每个线程完成N/(B*T)个向量的加法。
比如要计算33*1024个向量的加法,核函数调用为
add<<<128, 128>>>(a, b, c)
128*128个线程不能完成33*1024个向量的加法,核函数需要写成如下形式:
__global__ void add(int* a, int* b, int* c)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
while(tid < N)
{
c[tid] = a[tid] + b[tid];
tid += blockDim.x * GridDim.x;
}
}