CUDA经常使用概念及注意点

线程的索引计算

只须要知并行线程的初始索引，以及如何肯定递增的量值，咱们但愿每一个并行线程从不一样的索引开始，所以就须要对线程索引和线程块索引进行线性化，每一个线程的其实索引按照如下公式来计算：

int tid = threadIdx.x + blockIdx.x * blockDim.x;

线程块数量限制：65536

线程块中线程数量限制：512

共享内存和同步

共享内存

__share__添加到变量声明,使得声明的变量驻留在共享内存中。cuda c编译器对共享内存中的变量和普通变量采用不一样的处理策略，对于GPU启动的每一个线程块，cuda c都将建立该变量的一个副本。线程块中的全部线程都会共享这块内存。但线程却没法看到也不能修改其余线程块的变量副本。这是同一个线程块中的不一样线程进行通讯和协做的基础。

共享内存缓冲区驻留在物理gpu上，访问时延极低。

同步

__syncthreads();对线程块中的线程进行同步。线程发散的容易出现，使得部分场景下的线程同步颇有必要。

线程发散：当某些线程须要执行一些指令，而其余线程不须要执行时，这种状况叫作线程发散。在正常环境中，发散的分支会使得某些线程处于空闲状态，而其余线程将执行线程中的代码。在__syncthreads()状况中，线程发散形成的结果有些糟糕，cuda架构将确保，除非线程块中全部的线程都执行了同步操做，不然没有任何线程能够执行同步操做以后的指令。

常量内存与事件

常量内存：NVIDIA提供64k的常量内存，有效减小内存宽带。__constant__ 将变量的访问限制为只读。

从主机内存复制到GPU上的常量内存，使用方法cudaMemcpyToSymbol()进行复制。

性能提高缘由

（1）对常量内存的单次操做能够广播到其余邻近线程，节约15次的读写操做；

当处理常量内存时，NVIDIA硬件将单次内存读取操做广播到每一个半线程束。

（2）常量内存的数据将缓存起来，所以对相同地址的连续访问不会产生额外的内存通讯量。

线程束：warp

在cuda架构中，线程束指的是一个包含32个线程的集合，这些个线程被编制在一块儿，而且以步调一致（LockStep）的形式执行，在程序中的每一行，线程束中的每一个线程都将在不一样的数据上执行相同的命令。

事件API

cuda的事件本质上其实就是一个时间戳，这个时间戳就是在用户指定的时间上记录的。得到一个时间戳只有两个步骤：建立一个事件，记录一个事件。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
// gpu执行操做
...
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);

cudaEventDestroy(start);
cudaEventDestroy(stop);

文理内存（Texture Memory）

简介：

​ 与常量内存相似，只读，缓存在芯片上，减小对内存的请求，并提供更高效的内存带宽。专门为在内存访问模式中存在大量空间局部性（special locality）的图形应用程序而设计。在某个计算应用程序中，这意味着一个线程读取的位置可能与邻近线程读取的位置很是接近。纹理内存专门为加速这种内存访问模式。

纹理存储器中的数据是以一维、二维或者三维数组的形式存储在显存中，能够经过缓存加速访问，而且能够声明大小比常量存储器要大得多。在kernel中访问纹理存储器的操做称为纹理拾取。将显存中的数据和纹理参考系关联的操做，称为将数据和纹理绑定。显存中能够绑定到纹理的数据有两种，分别是普通的线性存储器和cuda数组。

使用步骤：

（1）须要将输入的数据声明为texture类型的引用；声明变量在gpu上；

（2） gpu中分配内存，经过cudaBindTexture()将变量绑定到内存缓冲区。告诉cuda运行时两件事情：

• 咱们但愿将指定的缓冲区做为纹理来使用；
• 咱们但愿将纹理引用做为纹理的“名字”。

（3）启动核函数，读取核函数中的纹理时，须要经过特殊的函数来告诉GPU将读取请求转发到纹理内存而不是标准的全局内存，使用编译器内置函数：tex1Dfetch();

（4）释放缓冲区，清除与纹理的绑定，cudaUnbindTexture();

流

页锁定内存

页锁定主机内存，固定内存，不可分页内存，OS将不会对这块内存分页而且交换到磁盘上。从而确保该内存始终驻留在物理内存中。OS能够安全的使某个应用程序访问该内存的物理地址，这块内存将不会被破坏或者从新定位。

• malloc分配的是标准的、可分页的主机内存；
• cudaHostAlloc将分配页锁定的主机内存。

建议：仅对cudaMemcpy（）调用的源内存或者目标内存，才能使用页锁定内存，而且在不须要使用他们时，当即释放。

流

支持设备重叠功能的设备，支持设备重叠功能的GPU可以在执行一个CUDA C核函数的同时，还能在设备和主机之间进行复制操做。

一些新的GPU设备同时支持核函数和两次的复制操做，一次是从主机到设备，一次是从设备到主机。在任何支持内存复制和核函数的执行相互重叠的设备上，当使用多个流时，应用程序的总体性能都能获得提高。

判断设备是否支持计算与内存复制操做的重叠：

int main( void ) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR( cudaGetDevice(&whichDevice) );
	HANDLE_ERROR( cudaGetDeviceProperties(&prop, whichDevice) );
	if(!prop.deviceOverlap) {
      	printf("Device will not handle overlaps");
      	return 0;
	}
}

多GPU系统上的CUDA C

零拷贝内存：能够在cuda C核函数中，直接访问这种类型的主机内存，因为这种内存不须要复制到GPU，所以称为零拷贝内存。经过cudaHostAlloc进行分配，最后一个参数采用：cudaHostAllocMapped.

判断设备是否支持映射主机内存：

int main( void ) {
	cudaDeviceProp prop;
	int whichDevice;
	HANDLE_ERROR( cudaGetDevice(&whichDevice) );
	HANDLE_ERROR( cudaGetDeviceProperties(&prop, whichDevice) );
	if(prop.canMapHostMemory != 1) {
      	printf("Device can not map memory");
      	return 0;
	}
}

当输入内存和输出内存都只是用一次时，那么在独立GPU上使用零拷贝内存将带来性能提高。

判断某个GPU时集成的仍是独立：

cudaGetDeviceProperties()获取属性结构体，该结构中的域：integrated，若是是设备是集成GPU，该值为true，不然为false。

注意：多GPU场景下，每一个gpu若是都要运行gpu程序的话，都须要主机cpu启动单独的线程进行资源控制，都有对应本身的线程。

《Programming Massively Parallel Processors: a Hands-On Approach》