CUDA并行编程较有用的总结

Cuda并行编程学习时候需注意的一些基本概念

1、Cuda的编程风格：spmp（单程序多数据）的并行编程风格。

2、在多GPU下，cudaMemcpy()不能用于GPU之间的数据复制

3、cudaMemcpy（）传输的数据类型有四种：

（1）       主机-主机

（2）       主机-设备

（3）       设备-主机

（4）       设备-设备

4、

（1）函数声明前面加_global_关键字的时候，表明这个函数是kernel函数并且从主机调用它的时候会在设备上生成线程网格。

（2）函数前加_device_关键字表明正在声明的函数是cuda设备函数，此函数只能在设备上执行，并且只能从kernel函数或者其它的设备函数调用。

（3）顾名思义_host_表明是函数是主机代码，一般不用加，因为大部分cuda程序都是从cpu代码移植过来的，默认就是主机函数

（4）_host_和_device_同时使用时触发编译系统，生成同一函数的两个不同的版本，它支持一种常见的应用，即只需要重编译同一函数的源代码就可以生成一个在设备上运行的版本。

5、在cuda应用中,所有线程同时并行的执行同一段kernel代码，线程在运行时访问硬件寄存器，硬件寄存器为线程提供可标识的坐标（通过线程索引threadIdx.x和threadIdx.y来区分不同的线程）

6、kernel函数中参数grid block thread之间的关系需要弄明白

一个网格grid是一个二维或者多维的数组，其中的每个元素代表一个Block

对一个GPU设备而言，一个block块包含有固定数目的thread（通成为512或者1024）

在执行kernel函数需要设置这些执行配置参数

例如

Dim3 dimBlock(width,width);//描述块的配置信息                 

Dim3 dimGrid(1,1);        //描述网格的配置信息

Kernel<<<dimGrid,dimBlock>>>(……..)  //启动在设备上进行计算的线程，给出了网格的维度（每个网格中块的数目）和块的维度（每一个块中的线程数目）

7、cuda的API函数为cuda程序提供服务，例如cudaMalloc()和cudaMemcpy()等等，可以在CUDA programming Guide中查看更多的API函数。

8、在cuda线程组织结构中，同一个网格中的线程执行同一个kernel函数，不同的网格间执行不同的kernel函数。

9、(1)线程最简单的组织结构就是把网格层次中的所有块都组织成一维数组的形式，每个块中的线程也组织成一维数组的形式。此时每个线程对应的变量（偏移地址）地址threadID=blockIdx.x*blockDim.x+threadIdx.x

(2)一般情况下我们把一个网格中的线程块组织成二维数组的形式，把每个块中的线程组织成三维数组的形式，具体的组织是由启动kernel时提供的执行配置决定。例如：

Dim3 dimGrid(2,2,2);

Dim3 dimBlock(4,2,2);

kernelFunction<<<dimGrid,dimBlock>>>(..........);

上面展示了一个简单的2*2的网格，并且每个块都由二维坐标组成，此时要标识一个线程的位置就要用到四个变量（blockIdx.x,blockIdx.y）和（threadIdx.x,threadIdx.y）。

10、cuda中用栅栏同步函数_syncthread()_来协调同一个块中的线程。Cuda块之间不需要同步，cuda运行时候系统可以任意相对顺序执行这些块，这种灵活性使得可扩展的实现称为可能。

11、warp线程调度和容许延时，在支持cuda的硬件实现方式中，一旦一个块分配到一个SM（多核流处理器），该块就被划分为32个线程的单元，这样的一个单元叫做warp.warp的大小是在具体实现时指定的，它并不是cuda规范中的一部分，但是知道warp方面的知识有助于我们理解和优化在特定cuda设备上执行cuda应用程序的性能。

在一个SM中，warp是线程调度的基本单元，每个warp包含32个线程，这些线程的threadIdx的值是连续的。比如有三个block被分配到一个SM中，而每一个block中又分为三个warp。

到底为什么要用warp呢，如果每个SM中只有8个SP，那么为何一个SM中需要这么多的warp呢？答案在于CUDA处理器需要高效的执行长延时操作，比如访问全局存储器。如果warp中的线程执行一条指令需要等待前面启动的长延时操作的结果，那么就不会选择执行该warp，而转向执行另一个不用等待结果的驻留的warp。在多个warp准备执行的时候就采用优先机制选择一个warp.这种机制让不产生延时的线程先执行，这就是所谓的延时隐藏（latency hiding）.只要有足够多的warp,任何时候硬件都可能找到适合的warp来执行，选择处于就绪状态的warp不会带来多余的时间开销，这就是所谓的“零开销线程调度（zero-overhead thread scheduling）”。GPU之所以具有这种允许这种长延时操作的能力，是因为它不同于CPU，CPU采用缓存和分支预测机制，从而牺牲了芯片的面积。

12、对于存储器的访问效率，它主要通过浮点计算与全局存储器访问操作之间的比值来衡量，这个比值又称为CGMA（Compute To Global Memory Access）,它的定义是在CUDA程序的某一区域内每次访问全局存储器时执行浮点运算的次数。

13、CUDA变量类型                 存储器              作用域                   生命周期

    除自动数组变量外的自动变量   寄存器              线程                       kernel函数

   自动数组变量                  局部存储器          线程                       kernel函数

    _device_ _shared_ int sharedVar 共享存储器       块                         kernel函数

 _device_ int GlobalVar           全局存储器         网格                       应用程序

_device_ _constant_  int ConstVar  常数存储器        网格                       应用程序

*********************************************************

_shared_ 提供同一块中线程的协同

_device_一个关键字表示变量是全局变量

_constant_变量存储在全局存储器中，但是采用缓存提高访问效率，并且采用合适的访问模式，速度快，可并行访问。

 

目前不同块中的线程还没办法实现同步，全局变量常用于调用一个kernel函数和另一个kernel函数时传递消息。

cuda变量在使用指针时有所限制，一般而言，指针用于指向全局存储器中的数据对象。在kernel和设备函数中，指针的使用方式有两种，第一：如果一个对象是由主机端的函数分配的，则指向对象的指针由cudaMalloc()初始化，并且作为参数传递给kernel函数。第二：全局存储器中变量的地址赋予指针变量，如float *ptr=&GlobalVar。

转载来自：http://blog.csdn.net/mysniper11/article/details/8270137