内建变量与线程索引

****执行配置运算符<<< >>>，用来传递内核函数的执行参数：
执行配置有四个参数，
第 一个参数声明网格的大小，
第二个参数声明块的大小，
第三个参数声明动态分配的 共享存储器大小，默认为 0，
最后一个参数声明执行的流，默认为 0。

****五个内建变量，用于在运行时获得网格和块的尺寸及线程索引等信息：
1. gridDim, gridDim 是一个包含三个元素 x,y,z 的结构体，分别表示网格在 x,y,z 三个方向上的尺寸，虽然其有三维，但是目前只能使用二维； 线程格的维度

2, blockDim, blockDim 也是一个包含三个元素 x,y,z 的结构体，分别表示块在 x,y,z 三个方向上的尺寸（blockDim.x，blockDim.y，blockDim.z）
对应于执行配置中的第一个参数，对应于执行配置 的第二个参数； 线程块的维度

3, blockIdx, blockIdx 也是一个包含三个元素 x,y,z 的结构体，分别表示当前线 程所在块在网格中 x,y,z 三个方向上的索引（blockIdx.x，blockIdx.y，blockIdx.z）； 线程块的ID索引（block在grid中的位置）

4, threadIdx, threadIdx 也是一个包含三个元素 x,y,z 的结构体，分别表示线程在其所在块中 x,y,z 三个方向上的索引； 线程ID的索引（即线程在block中的位置）

5, warpSize，warpSize 表明 warp 的尺寸，在计算能力为 1.0 的设备中，这个值是 24，在 1.0 以上的设备中，这个值是 32。
6. 对于一维的block，线程的 threadID
= threadIdx.x



7. 对于大小为（blockDim.x,
blockDim.y）的 二维 block，线程的 threadID=threadIdx.x+threadIdx.y*blockDim.x



比如thread(1,1)的相对位置为： 1+ 1*16=17

8. 对于大小为（blockDim.x, blockDim.y, blockDim.z）的 三维 block，线程的 threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y



实例：

__global__ void matrixAddition(float *a, float *b, float *c)
{
int x = blockIdx.x * blockDim.x + threadIdx.x ;
int y = blockIdx.y * blockDim.y + threadIdx.y ;
int dimGx = gridDim.x;
int dimGy = gridDim.y;
int offset = y * dimGx * blockDim.x + x;
c[offset] = a[offset] + b[offset];
}

// 4个Block ,每个Block为20*20个thread
// float a[N * N], b[N * N], c[N * N];//40*40
int gridSize = 2*2;
dim3 blockSize(N/2, N/2);//20 * 20
matrixAddition<<<gridSize, blockSize>>>(dev_a, dev_b, dev_c);