cuda 线程块与线程

一维网格

    假设一个数组长度N为64，我们设置线程数TPB=32,则线程块数为N/TPB=2，核函数可写为：

Kernel<<<N/TPB,TPB>>>()

    如果设置N=65，按照N/TPB我们仍然会得到2个包含32个线程的线程块，数组的最后一个输入没有被覆盖到，有一个简单的办法可以保证线程网格覆盖到整个数组程度，就是声明线程块数目从N/TPB变为(N+TPB-1)/TPB:

Kernel<<<(N+TPB-1)/TPB,TPB>>>()

    但此时网格的索引可能超出数组范围并引发段冲突错误，我们需要在核函数中加入一个控制语句来防止此错误。
    注意线程网格和线程块的维度实际是三维的类型dim3，未指明的分量默认设为1。执行配置<<<len/TPB,TPB>>>将被解释为<<<dim3(len/TOB,1,1),dim3(TPB,1,1)>>>

二维网格

    在CUDA中从一维到二维过渡仅需很少的调整就可以。在一维中，我们指定线程块和网格的大小为整数，计算出索引值 i:

int i = blockIdx.x*blockDim.x + threadIdx.x;

在二维中引入行索引和列索引，计算如下：

int c = blockIdx.x*blockDim.x + threadIdx.x;
int r = blockIdx.y*blockDim.y + threadIdx.y;

我们在一维数组中以行为主序来存储矩阵元素，图像宽为w，因此一维数组中索引i的计算如下：

int i = r*w + c;

我们用一个声明来指定二维线程块的大小：

dim3 blockSize(TX,TY);//Equivalent to dim3 blockSize(TX,TY,1);

然后我们像一维那样计算x、y方向上的线程块数(bx和by)：

int bx = (W + blockSize.x - 1) / blockSize.x;
int by = (H + blockSize.y - 1) / blockSize.y;

指定网格大小(在线程块中)的语法是:

dim3 gridSize = dim3(bx,by);

启动内核程序：

kernelName<<<dSize,blockSize>>>(args)