cuda教程

注意区分X轴和Y轴

                  研究CUDA已经有段时间了，但是只是皮毛罢了，糊弄糊弄外行人还可以，但真是情况是。。。此处省略几个字，遂回头整理重温一下，所谓温故而知新。

在上图中, blockIdx.x = 0, blockDim.x = 10 ， threadIdx.x从0变化到9.

在内核函数内部，CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：

      1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，
它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1;

      2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1;

      3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸。拿上图来说，注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4, blockDim.y = 3;

      4. threadIdx：代表线程的索引，如Thread(3,2)，则threadIdx.x=3,threadIdx.y=2;

  每个CUDA线程拥有独立的本地内存（local Memory）；每一个线程块（block）都有其独立的共享内存（shared memory），共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存（global memory）的区域对所有的CUDA线程都是可访问的。
        除了上述三种存储资源以外，CUDA还提供了两种只读内存空间：常量内存（constant memory）和纹理内存（texture memory），同全局内存类似，所有的CUDA线程都可以访问它们。对于一些特殊格式的数据，纹理内存提供多种寻址模式以及数据过滤方法来操作内存。这两类存储资源主要用于一些特殊的内存使用场合。

block是有维度的，一维、二维、三维。

对于一维的block，tid = threadIdx.x

当block和grid都是一维的时候，tid=threadIdx.x+bolckDim.x*blockIdx.x;

对于（Dx，Dy）二维的block，tid = threadIdx.x + blockIdx.x*blockDim.x    //（blockDim.x线程块的大小即线程的数量，blockIdx.x为线程块的索引）

我们将具体点的，在主机函数中如果我们分配的是这样的一个东西：

dim3 blocks(32,32);

dim3 threads(16,16);（与dim3 threads（16,16,1）等价）

dim3是神马？dim3是一个内置的结构体，dim3结构变量有x，y，z，表示3维的维度,如果只需要二维，定义时只需将第三元缺省或置为1。

kernelfun<<<blocks, threads>>>();

我们调用kernelfun这个内核函数，将blocks和threads传到<<<,>>>里去，这句话可牛逼大了——相当于发号施令，命令那些线程去干活。这里使用了32*32 * 16*16个线程来干活。你看明白了吗？blocks表示Gird里有二维的32*32个block，而每个block中又用了16*16的二维的thread组。好吧，我们这个施令动用了262144个线程！

tips：如果矩阵的尺寸不能背block的尺寸整除，则grid的尺寸应该为dim3 sizeofGrid（NUM_X / blocks_x+1，NUM_Y / blocks_y+1）;NUM为矩阵的尺寸.

_syncthreads()

，可以在内核中调用__syncthreads()内置函数指明同步点；__syncthreads()起栅栏的作用，在其调用点，块内线程必须等待，直到所以线程都到达此点才能向前执行