CUDA 学习笔记 3

CUDA 编程模型

回顾前面的例子：

__global__ void addKernel(int *c, const int *a, const int *b){    int i = threadIdx.x;    c[i] = a[i] + b[i];}int main(){    ...    addKernel <<< 1, size >>>(dev_c, dev_a, dev_b);    cudaStatus = cudaDeviceSynchronize();    ...    return 0;}

kernel代码是在GPU上运行的，而主函数中kernel代码前后很多准备和善后工作是在CPU上执行的，这种由多种不同结构的处理器协作运行应用的运行模式称为异构计算模型。
CUDA的编程模型就是异构的。在这种模型下，主CPU称为主机（Host），GPU（当然也可以是其他协处理器）称为设备（Device），Host和Device各自维护自己独立的DRAM内存，也分别有各自的内存空间，Host的内存就是系统内存，GPU的内存就是设备内存，也就是GPU的显存。Host和GPU不能任意访问对方的内存空间（CUDA6 开始支持Unified Memory模型，这是后话，暂且不表），必须通过PCIe总线访问，在应用程序中，就需要调用CUDA Runtime，通过CUDA 驱动程序来读写GPU显存，以及分配或释放这些空间。

CUDA线程模型

在调用GPU并行执行 kernel 的时候，使用 <<< a, b >>> 来指定所需的线程数量。其中 a 是线程块(Block)的数量，b 是每个Block中的线程数量。但事情比想象的要略复杂一些，因为这里 a 和 b 可以是int 型的整数，也可以是dim3 型的向量。dim3 是CUDA定义的数据类型，是一个不超过3维的向量，比如：

dim3 v1(10);            // v1 是1维向量dim3 v2(10, 10);        // v2 是2维向量dim3 v3(10, 10, 10);    // v3 是3维向量

也就是说可能有 (x * y * z) 个Block，每个 Block里有可能有 （u * v * w）个线程在执行。CUDA的线程就是通过这种方式组织起来的，可以用下图来直观的说明（图中每个人代表一个线程）：

1. 假设任务简单，只要一个班组（假设3个人）并行处理就可以了:
   
   那相应的线程配置为：<<< 1, 3 >>>

2. 如果任务稍微复杂一点，那就再安排一层楼的所有班组（假设共3组，每组3人）：
   
   那相应的线程配置为：<<< 1, (3,3) >>>

3. 如果任务更复杂，需要整栋厂房所有楼层的所有班组（假设共2层楼）并行工作：
   
   那相应的线程配置为：<<<1, (4, 3, 2)>>>

至此，一栋厂房（CUDA中的Block）的工人都派去工作了。假设任务再复杂，那就不仅是一栋厂房的事了，需要协调其他厂房，这就要修改厂房数量（Block 数）：
<<< 5, (4, 3, 2)>>> // 整个工厂中的第一条马路上的5个厂房
<<< (5, 6), (4, 3, 2)>>> // 整个工厂中的6条马路，每条马路上的5个厂房
<<< (5, 6，7), (4, 3, 2)>>> // 整个公司中的7个分工厂，每个厂的6条马路，每条马路上5个厂房

这样就把一个公司所有的工人全调动起来了。
由一个单独的kernel启动的所有线程组成的整个线程组称为Grid（相当于上例中的整个公司），所以CUDA的线程分组包含Grid和Block 两个层次，<<<a, b>>> 中的a 就称为 Grid维度，b称为 Block维度。下面是一个2维Grid，2维Block组成的线程Grid示意图：

Grid中的每一个线程都会有一个线程Id，CUDA内置了几各变量（不是关键字），可以用来访问线程Id：

变量 含义 gridDim grid维度，就是前面说的a blockDim block维度，就是前面说的b threadIdx 线程的index blockIdx block的index

如果grid的维度是dim3向量，那么可以通过gridDim.x, gridDim.y分别得到x方向和y方向的维度。比如上图中gridDim = (3，2)，即gridDim.x = 3, gridDim.y=2。
这时blockIdx就是2维向量，当blockIdx=（1，1）的时候，就是访问上图下方的那个Block，同理，这个block里的线程的threadIdx也是二维的，比如第一个线程的threadIdx.x=0,threadIdx.y =0。
由此，一个block里线程的Id可以通过threadIdx计算出，分为如下三种情况：

Block维度 线程Id 1 等于 threadIdx 2 threadIdx.x + threadIdx.y * blockDim.x 3 threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y

CUDA之所以这样组织线程，是因为GPU处理的图形数据往往就是2维或者3维的，所以硬件的线程就是按类似的逻辑组织的，而且在大规模数据计算的时候，遇到的数据也往往就是2维或者3维的，那么自然地，这样组织线程比较直观也比较有效率。

但是线程配置仍然不是任意的：每个block 中线程总数不能超过1024，即blockDim.x * blockDim.y * blockDim.z <= 1024 。

之所以有这个限制，那又得从GPU硬件架构，以及内存模型说起。

CUDA硬件架构

下图是NVIDIA Kepler GPU 硬件架构：

其中有15个SMX（不同型号的GPU，SMX数量不一定相同），这个SMX 就是Stream Multiprocessors的意思，每个SMX里面有192个单精度Core和64个双精度Core，4个线程调度器Warp Scheduler。

CUDA内存模型

CUDA程序编译