CUDA编程经验技术总结 系列之《核函数之线程技术》

       关于GPU编程的这些资料均是我早期的一些资料，趁出差这段时间整理下，所以就直接复制过来了，其中会有一些瑕疵，请读者朋友斧正，如果有时间，接下来我会编写并分享使用gpu编程实际应用过程中的经验教训和总结。



0、核函数调用方式：kernel <<< a,b>>>();

gridDim俗名线程格，是二维的（第三维为1），而blockDim是线程块，是三维的，cuda运行时容许启动一个二维线程格，并且线程格中的每个线程块都是一个三维的线程数组（只要你愿意）。

一、N个一维线程块并行，每个线程块开启一个线程:

核函数：   kernel<<< N,1>>>();

线程索引： int  tid=blockIdx.x; 

其中blockIdx.x =1~N,   blockIdx.y=0;

threadIdx.x=1,    threadIdx.y=0;

二、一个一维线程块，该线程块开启N个一维并行线程:

核函数：   kernel<<< 1,N>>>();

线程索引： int  tid=threadIdx.x;

其中blockIdx.x =1,       blockIdx.y=0;

threadIdx.x=1~N,    threadIdx.y=0;

三、N个一维线程块并行，每个线程块开启M个一维并行线程:

核函数：   kernel<<< N，M>>>();

线程索引： int   tid=threadIdx.x+blockIdx.x*blockDim.x;(二维索引空间转换为线性空间的标准算法)

tid  += blockDim.x*gridDim.X；

其中  blockIdx.x=1~N,   blockIdx.y=0;

threadIdx.x=1~M,    threadIdx.y=0;

 注意：N，M不能超过1536或者1024，具体多少记不清了

blockDim对于所有线程块来说，是一个常量，保存的是线程块中每一维的线程数量。由于使用的是一维线程块，因此blockDim.y等于1，只使用blockDim.x=M就ok。

四、二维线程块并行，每个线程块开启M个一维并行线程:

Dim3 dev_block(a,b);

核函数：   kernel<<< dev_block，M>>>();

线程索引： int   tid=threadIdx.x+blockIdx.x*blockDim.x;(二维索引空间转换为线性空间的标准算法)

其中 blockIdx.x =0~a-1,   blockIdx.y=0~b-1;

threadIdx.x=1~M,    threadIdx.y=0;

blockDim对于所有线程块来说，是一个常量，保存的是线程块中每一维的线程数量。由于使用的是一维线程块，因此blockDim.x=M，blockDim.y=1。

gridDim.x=a,  gridDim.y=b;

五、N个一维线程块并行，每个线程块开启一个二维并行线程:

Dim3  dev_threads(a,b);

核函数：   kernel<<< N，dev_threads >>>();

线程索引： int   tid=threadIdx.x+blockIdx.x*blockDim.x;(二维索引空间转换为线性空间的标准算法)

其中 blockIdx.x =1~N,   blockIdx.y=0;

threadIdx.x=0~a,    threadIdx.y=0~b;

blockDim对于所有线程块来说，是一个常量，保存的是线程块中每一维的线程数量。由于使用的是一维线程块，因此blockDim.y=b，blockDim.x=a。

gridDim.x=1,  gridDim.y=0;

六、二维线程块并行，每个线程块开启二维并行线程:

Dim3 dev_block(a,b);

Dim3 dev_threads(c,d);

核函数：   kernel<<< dev_block，dev_threads >>>();

线程索引： int   tid=threadIdx.x+blockIdx.x*blockDim.x;(二维索引空间转换为线性空间的标准算法)

其中 blockIdx.x =0~c,   blockIdx.y=0~d;  threadIdx.x=1~a,  threadIdx.y=0~b;

blockDim对于所有线程块来说，是一个常量，保存的是线程块中每一维的线程数量。由于使用的是一维线程块，因此blockDim.y=d，blockDim.x=c。

gridDim.x=a,  gridDim.y=b;