浅谈GPU编程.ppt

《浅谈GPU编程.ppt》由会员上传分享，免费在线阅读，更多相关内容在行业资料-天天文库。

1、CUDA略谈引言逻辑线程拓扑CUDA存储器类型设备端存储器主机端存储器参考书CUDA统一计算设备架构它是一种将GPU作为数据并行计算设备的软硬件体系。开发GPU通用计算程序，需掌握并行算法和GPU架构方面的知识。支持CUDA的GPU可有效利用原用于图形渲染的计算资源进行通用计算。此处就基于GPU的CUDA编程谈一些认识。线程逻辑拓扑结构线程逻辑拓扑结构线程拓扑分两个层次。顶层是2维网格平面，以块为单位。块为3维立方结构，以线程为单位，故称线程块。grid:4×4block:4×4×4threadblockCUDA程序结构内核函数：即设备端代码。形如：__global__kernel(形

2、参){┋}主机端代码voidmain(){┋dim3block(4,4,4);dim3grid(4,4,1);kernel<<>>(实参);┋}存储器层次结构存储器类型设备端存储器结构寄存器共享存储器局部存储器常数存储器全局存储器纹理存储器板载显存GPU片上寄存器GPU片上高速存储器每个寄存器32bit。每个SM有大量寄存器，但由块内线程共享，故平均到每个线程的寄存器就很有限了。在内核函数中声明的少量变量是寄存器变量，每个线程都维护各自的寄存器变量，它们是线程私有的。__global__kernel(){intbx=blockIdx.x;intby=blockI

3、dx.y;inttx=threadIdx.x;intty=threadIdx.y;inttz=threadIdx.z;┋}局部存储器板载显存如果每个线程使用了过多的寄存器，或声明了大型结构体或数组，或者编译器无法确定数组的大小，线程的私有变量就有可能会被分配到局部存储器中。例见P46共享存储器（SM）线程块内共享片上高速存储器静态分配动态分配__global__kernel(形参){__shared__intsm_static[16];extern__shared__intsm_dynamic[];┋}voidmain(){┋intsm_dynamic_size=32;kernel<<

4、>>(实参);┋}共享存储器（SM）SM的组织方式在1.x计算能力的设备中，sm共16KB，被划分为16个bank，每个bank的宽度为32位。32bitbank001632bitbankii16+i32bitbank15153132bitbank1117…………共享存储器（SM）SM访问的bank冲突：每个warp为32个线程，一个warp对sm的访问被分成两个half-warp的访问，只有同一个half-warp内的线程才可能发生bank冲突。前half-warp的线程与后half-warp的线程之间则不会发生bank冲突。共

5、享存储器（SM）bankhalf-warpbankhalf-warp顺序访问随机访问无bank冲突bankhalf-warp3间隔访问共享存储器（SM）bankhalf-warpbankhalf-warp2间隔访问8间隔访问2路bank冲突bankhalf-warp广播机制8路bank冲突无bank冲突常数存储器只读存储器，数据位于显存，但拥有缓存加速。空间较小，只有64KB。每个SM拥有8KB的常数存储器缓存。在函数外定义，作用范围为文件域。主机端和设备端函数均可见。使用方法：方法1：定义时初始化，内核函数中直接使用。方法2：先定义，后在主机端用函数赋值。常数存储器__constan

6、t__charp_HelloCUDA[11];//定义__constant__intt_HelloCUDA[11]={0,1,2,3,4,5,6,7,8,9,10};//定义并初始化__constant__intnum=11;//定义并初始化__global__staticvoidHelloCUDA(char*result){inti=0;for(i=0;i

7、CUDA,helloCUDA,sizeof(char)*11);┋}直接使用函数赋值全局存储器位于板载显存，占据显存的绝大部分，没有缓存。也称线性内存。可定义两种数据结构：线性存储器：用于存放主机端传过来的数据或存放将要回传给主机的数据。CUDA数组：用于纹理绑定。为有效利用带宽，必须遵循合并访问要求，并避免分区冲突。全局存储器分区(partition)冲突：在中高端GPU中一般有多个存储器控制器。每个存储器控制器对应的地址空间称为一个分区，连