CUDA编程注意事项.docx

CUDA编程注意事项1. kernel 以线程网格grid的形势组织线程。现在线程网格（grid）只能有一个。2.每个线程网格又由线程块（block）组成。线程块可以以三维方式组织变量dim3 线程块是GPU的执行单位，每个block中线程数目不能超过512个，block里面的线程按乱序执行，并不同步，在同一个时刻同一个block的线程并不执行相同的指令，可以用__syncthreads()来实现栅格同步。每个block可以通过shared memory来共享数据GPU硬件架构3. 第一级tcp包含2-3个sm流多处理器。4.第二级sm包括8个sp流处理器。5.第三级sp流处理器硬件的最小单位。所以每个sm中有16个基本执行单位6.一个sm中可以分配至多8个获得的线程块，交替执行来隐藏延时。7. warp是实际的执行单位，block被划分为更小的线程束，划分由设备的计算能力决定，sm中每发送一个指令sm中的sp会将指令执行四次，所以warp被设计为32个线程一组。每个warp中线程执行同一条指令，warp是软件上真正的执行单位。在考虑存储器方面应用更多的是half-warp，warp的数目之和在1.0和1.1 的设备上不能超过24个，在1.2和1.3的设备上不能超过32个。8.在执行的过程中如果要控制单个线程的行为，必须使用分支，这会大大降低效率，在开发的过程中尽量做到warp里面不分支。9. Grid 的x，y最大可以到65536，但是实际划分中x*y最大65536，即每个Grid中最大有65536个线程块10. block中的x，y的最大值为512，y的最大值为4但是三者之积必须小于768（1.0和1.1）或1024（1.3和1.2）11. 最简单的__global__程序a)分配shared memoryb)将global memory的数据读入shared memory中c)进行计算并将结果放回到sharedmemory中d)将sharedmemory中的数据放到globalmemory中在理想的状态下在所有的存储器传输的同时GPU的各个核心都在进行计算这样能够很好的隐藏延时12. extern __shared__ memoryextern在设备端表示分配动态数组，在主机端表示外部变量声明。如果没有extern 则表示在sharedmemory中静态分配数组13. CUDA内建变量a) gridDim用于描述网格被划分的情况gridDim.xgridDim.y一个grid只有二维b) blockIdxblock的index blockIdx.x描述block在当前的grid中的位置取值范围[0,gridDim.x-1]blockIdx.yc) blockDimd)threadIdxe) warpSize1. CUDA计时函数cutCreateTimer(int)cutStartTimer(int) cutStopTimer(int)cutGetTimer(int)cutDeleteTimer(int)线程间的通信1同步函数保证一个block里的所有函数同步执行__syncthreads（）但是使用时要保证所有的thread都能执行到这一步，不然会造成死锁2. memory fence函数a) __threadfence(); 对grid中的所有线程可见b) __threadfence_block();对block中的所有现场可见c)二者的使用范围不同，但都是通知其他线程自己的操作已经完毕3. kernel间的通信，通过global memory4.多个GPU之间通信a)注意通过主机的内存通信，mapped memory 允许多个GPU访问同一个pinned memory5. GPU和CPU之间的通信a) CudaThreadSynchronize()保证cpu和Gpu同步执行异步并行执行6. kernel的启动7.以Async为后缀的内存拷贝函数8. device到device的内存拷贝函数9.存储器初始化函数10.流程序是通过流来管理并发的，每个流是按照顺序执行的一系列操作，而不同的流之间的乱序可能是乱序，也可能是并序a)流的创建cudaStreamCreate(cudaStream_t)b)流的销毁cudaStreamDestroy()c)流的查询cudaStreamQuery()d)流的等待cudaStreamSynchronize();e)线程的等待cudaThreadSynchronize();