CUDA笔记2：概念理解
CUDA全称是ComputeUnified Device Architecture，中文名称即统一计算设备架构，它是NVIDIA公司提出了一种通用的并行计算平台和编程模型。使用CUDA，我们可以开发出同时在CPU和GPU上运行的通用计算程序，更加高效地利用现有硬件进行计算。并行编程的中心思想是分而治之：将大问题划分为一些小问题，再把这些小问题交给相应的处理单元并行地进行处理。在CUDA中，这一思想便体现在Grid, Block, Thread等层次划分上。
GPU并行架构：
GPU编程中CPU被称为Host, GPU被称为Device.
Kernel函数在Host中被调用，在Device中被执行。
一个Kernel对应一个Grid;
一个Grid包含一组Block,Block在Grid中的分布可以是一维，二维或三维的，具体由GridDim定义，每个Block都有各自的ID,记blockIdx.xyz;
一个Block包含一组Thread,Thread在Block中的分布有BlockDim定义，每个Thread都有各自的ID，记threadIdx.xyz;
Kernel函数调用方式为kernel<<<A，B>>>(parameters)，在尖括号中，A代表线程格（grid）的尺寸gridDim，它可以是三维的，用类型dim3表示，也可以是一维的，用int类型表示。B代表线程块（block）的尺寸blockDimk，它与A类似，也可分别用dim3或int类型表示。在内核函数内部，CUDA为我们内建了一些变量用于访问Grid、Block的尺寸和索引等信息，它们是：
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：线程索引，同block索引。
从GPU硬件角度，它的基本单元为StreamingProcessor(SP),即流水处理器，每个SP执行一个Thread;多个SP组成一个Streamming MultiProcessor(SM),即流水多处理器，每个SM执行一个Block.GPU通常包含多个SM。
GPU存储层次：
每个 thread都有自己的一份 register 和 local memory的空间。一组thread构成一个 block，这些thread则共享有一份shared memory。__syncthreads()可以同步一个Block内的所有线程，不同Block内的Thread不能同步。此外，所有的 thread(包括不同 block 的 thread)都共享一份global memory、constant memory、和 texture memory。不同的 grid则有各自的 global memory、constant memory和 texture memory。cudaMalloc和cudaFree用于内存分配及释放，它们分配的是global memory，cudaMemcpy用于Hose-Device数据交换。
一个Thread被执行过程：
Grid在GPU上启动；
block被分配到SM上；
SM把线程组织为warp；
SM调度执行warp；
执行结束后释放资源；
block继续被分配....
例如GTX760, 6SM， 192SP/SM，一个SM一次执行一个Block，假设一个Warp含32Thread,一个Block线程数量应该远远大于192(6warp)，为的是GPU执行长延时操作。（CUDA处理器需要高效地执行长延时操作，如果warp中的线程执行一个条指令需要等待前面启动的长延时操作的结果，那么不会选择执行该warp，而是选择执行另一个不用等待结果的驻留的warp，这样，如果有了多个warp准备执行，则总可以选择不产生延时的线程先执行，达到所谓的延时隐藏。）
函数类型：
- __device__ // 执行于Device，仅能从Device调用。限制，不能用&取地址；不支持递归；不支持static variable；不支持可变长度参数
- __global__ // void： 执行于Device，仅能从Host调用。此类函数必须返回void
- __host__ // 执行于Host，仅能从Host调用，是函数的默认类型
- // 在执行kernel函数时，必须提供execution configuration，即<<<....>>>的部分。
- // 例如：
- __global__ void KernelFunc(...);
- dim3 DimGrid(100, 50); // 5000 thread blocks
- dim3 DimBlock(4, 8, 8); // 256 threads per block
- size_t SharedMemBytes = 64; // 64 bytes of shared memory
- KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>>(...);
数据结构：
- // 内建矢量类型：
- int1，int2，int3，int4，float1，float2， float3，float4 ...
- // 纹理类型：
- texture<Type, Dim, ReadMode>texRef;
- // 内建dim3类型：定义grid和block的组织方法。例如：
- dim3 dimGrid(2, 2);
- dim3 dimBlock(4, 2, 2);
- // CUDA函数CPU端调用方法
- kernelFoo<<<dimGrid, dimBlock>>>(argument);
- __device__ // GPU的global memory空间，grid中所有线程可访问
- __constant__ // GPU的constant memory空间，grid中所有线程可访问
- __shared__ // GPU上的thread block空间，block中所有线程可访问
- local // 位于SM内，仅本thread可访问
- // 在编程中，可以在变量名前面加上这些前缀以区分。
CUDA计时：
CPU代码和CUDA函数是异步执行的，即CUDA函数完成任务之前控制权有可能已经交给了CPU. cudaDeviceSynchronize()能够阻塞当前CPU线程，直到CUDAkernel函数调用运行完毕后才继续执行CPU线程。所以如果需要对Host和Device代码进行联合计时，通常在CUDA Kernel函数和结果拷贝cudaMemcpy函数之间增加一句同步代码cudaDeviceSynchronize()。除了在CPU主程序中设置定时函数外，CUDAAPI也提过了毫秒级的GPU计时方法：
- cudaEvent_t start, stop;
- float time;
- cudaEventCreate(&start);
- cudaEventCreate(&stop);
- cudaEventRecord(start, 0);
- kernel<<<grid, threads>>>(d_odata, d_idata, size_x, size_y, NUM_REPS);
- cudaEventRecord(stop, 0);
- cudaEventSynchronize(stop);
- cudaEventElapsedTime(&time, start, stop);
- cudaEventDestroy(start);
- cudaEventDestroy(stop);
这里cudaEventRecord()用于将start和stop事件放置到默认流（stream0），当设备运行到事件处时，将会为该事件记录一个时间戳（timestamp）。cudaEventElapsedTime()函数返回记录中start和stop事件之间的毫秒级时间差。由于记录是发生在GPU上的，其时钟是独立的，因此时钟分辨率也是操作系统无关的
CUDA代码调试：
CUDA Toolkit安装时会自动安装VS插件Nsight。打开VS,在device code上设置断点，VS菜单栏选择Nsight-> Start CUDA Debugging，即可调试CUDA代码。
参考：
CUDA学习笔记：http://luofl1992.is-programmer.com/posts/38830.html
CUDA编程指南阅读笔记:http://blog.csdn.net/csgxy123/article/details/9704461
大规模并行处理器编程实战笔记：http://blog.csdn.net/linyingzhan/article/details/8265088