CUDA编程模型(入门).ppt

《CUDA编程模型(入门).ppt》由会员分享，可在线阅读，更多相关《CUDA编程模型(入门).ppt（25页珍藏版）》请在淘文阁 - 分享文档赚钱的网站上搜索。

1、本课程得到Intel大学合作计划支持课程网站：http:/ CUDA编程模型华南理工大学 陈虎 博士GPU与CPU的差异GPU(Graphics Process Unit)面向计算密集型和大量数据并行化的计算大量的晶体管用于计算单元通用通用CPU面向通用计算大量的晶体管用于Cache和控制电路DRAMCacheALUControlALUALUALUDRAMCPUGPUGPU与CPU的峰值速度比较1 Based on slide 7 of S.Green,“GPU Physics,”SIGGRAPH 2007 GPGPU Course.http:/www.gpgpu.org/s2007/slid

2、es/15-GPGPU-physics.pdf第一代GPU结构(GeForce 6800)第二代GPU(GeForece 8800)GeForce 8800的主要技术参数晶体管数目（百万）晶体管数目（百万）681工艺90nm芯片面积(mm2)470工作主频（GHZ）1.5峰值速度（Gflops）576处理器数目128片 上 存 储 器 容 量（KB）488功耗（W）15016个流多处理器(SM)每个SM中包含了8个流处理器SP每个SP包含一个乘加单元每个SM管理了24个线程簇(warp)，共有768个线程采用单线程多数据(Single-thread,Multiple-data)技术每个周期在8

3、个SP上并行执行一个线程簇GF8800的存储层次结构层次层次容量容量(KB)延迟延迟(ns)局部存储器1626L1 Cache5280L2 Cache32370DDR510GeForce 8800GPU最适合做什么？对多个数据进行同一种运算（STMD适用）一次存储器访问，多次运算（外部DDR访问开销高，局部存储器容量较小）浮点计算比例高（特别是单精度浮点）典型计算：物理模拟，线性代数计算应用领域：计算生物学图像处理CUDA编程模型Nvidia公司开发的编程模型可以和VC 8.0集成使用下载地址：http:/ to TargetCompiler G80 GPU Target codePTX Co

4、deVirtualPhysicalCPU Codefloat4 me=gxgtid;me.x+=me.y*me.z;ld.global.v4.f32$f1,$f3,$f5,$f7,$r9+0;mad.f32$f1,$f5,$f3,$f1;链接时的动态库：The CUDA runtime library(cudart)The CUDA core library(cuda)CUDA中的线程线程：CUDA中的基本执行单元；硬件支持，开销很小；所有线程执行相同的代码（STMD）每个线程有一个唯一的标识IDthreadIdx若干线程还可以组成块(Block)线程块可以呈一维、二维或者三维结构Device

5、Grid 1Block(0,0)Block(1,0)Block(2,0)Block(0,1)Block(1,1)Block(2,1)Block(1,1)Thread(0,1)Thread(1,1)Thread(2,1)Thread(3,1)Thread(4,1)Thread(0,2)Thread(1,2)Thread(2,2)Thread(3,2)Thread(4,2)Thread(0,0)Thread(1,0)Thread(2,0)Thread(3,0)Thread(4,0)CUDA存储器分配cudaMalloc()在全局存储器中分配空间两个参数：地址指针空间大小cudaFree()在全局存

6、储器中回收空间参数回收空间地址指针GridGlobalMemoryBlock(0,0)Shared MemoryThread(0,0)RegistersThread(1,0)RegistersBlock(1,0)Shared MemoryThread(0,0)RegistersThread(1,0)RegistersHost示范代码分配64*64的单精度浮点数组存储器地址为Md.elementsTILE_WIDTH=64;Matrix Md/*“d”is often used to indicate a device data structure*/int size=TILE_WIDTH*TI

7、LE_WIDTH*sizeof(float);cudaMalloc(void*)&Md.elements,size);cudaFree(Md.elements);主机和设备之间的数据传输cudaMemcpy()存储器数据传输参数：目的地址源地址传输字节数传输类型主机主机主机全局存储器全局存储器全局存储器异步传输GridGlobalMemoryBlock(0,0)Shared MemoryThread(0,0)RegistersThread(1,0)RegistersBlock(1,0)Shared MemoryThread(0,0)RegistersThread(1,0)RegistersHo

8、st示范代码传输 64*64 单精度浮点数组M在主机存储器中Md在设备存储器中传输方向常数：cudaMemcpyHostToDevicecudaMemcpyDeviceToHostcudaMemcpy(Md.elements,M.elements,size,cudaMemcpyHostToDevice);cudaMemcpy(M.elements,Md.elements,size,cudaMemcpyDeviceToHost);CUDA的函数声明在哪里执行:只能从哪里调用:_device_ float DeviceFunc()设备设备_global_ void KernelFunc()设备主机

9、_host_ float HostFunc()主机主机 _global_定义一个内核函数必须返回 void _device_ 和 _host_可以一起使用CUDA 的函数声明(续)_device_ 函数不能使用函数指针；_device_函数在设备上执行，所以:不能递归不能在函数内定义静态变量不能使用变量作为参数调用一个内核函数 线程生成内核函数必须使用执行配置调用:_global_ void KernelFunc(.);dim3 DimGrid(100,50);/5000 thread blocks dim3 DimBlock(4,8,8);/256 threads per block siz

10、e_t SharedMemBytes=64;/64 bytes of shared memoryKernelFunc(.);/用指明该函数调用的线程数任何从CUDA 1.0调用的内核函数都是异步的,阻塞需要显式的同步。CUDA变量类型修饰符当使用_local_,_shared_,和 _constant_时，_device_ 是可选的没有任何修饰符的自动变量将默认放在寄存器内数组除外，它们将会被摆放在本地存储器内变量声明存储器作用域生命期_device_ _local_ int LocalVar;本地存储器threadthread_device_ _shared_ int SharedVar;共

11、享存储器blockblock_device_ int GlobalVar;全局存储器gridapplication_device_ _constant_ int ConstantVar;常存储器gridapplication20在哪里声明变量?主机能否访问?在任何函数外在内核可以不可以全局存储器常存储器寄存器(自动)共享存储器本地存储器21G80 的CUDA存储器实现每一个thread 可以:可以读写寄存器可以读写本地存储器每一个block可以读写 共享存储器每一个grid 可以读写 全局存储器每一个grid 只可读 常数存储器GridGlobal MemoryBlock(0,0)Shared

12、 MemoryThread(0,0)RegistersThread(1,0)RegistersBlock(1,0)Shared MemoryThread(0,0)RegistersThread(1,0)RegistersHostConstant Memory22通用的编程策略全局存储器在设备内存中(DRAM)它的访问速度比共享存储器要慢在设备上执行运算的优化方法是：把数据分割成子集，放入共享存储器中以一个线程块操作一个数据子集：从全局存储器读入数据子集到共享存储器，使用多个线程实现存储级并行在共享存储器中对数据子集执行运算;每一个线程可以有效率地多次访问任何数据元素从共享存储器上拷贝结果到全局

13、存储器通用的编程策略(续)常数存储器也在设备内存里(DRAM)比访问共享存储器慢得多但是它可以高速缓存，对于只读数据的访问非常高效。通过访问模式小心划分数据只读 常数存储器(若在高速缓存内，非常快)在块内读写共享 共享存储器(非常快)线程可读写 寄存器(非常快)读写输入和结果 全局存储器(非常慢)24变量类型约束指针只能指向在全局存储器已分配或声明的内存:在主机分配，然后传递给内核:_global_ void KernelFunc(float*ptr)以全局变量的地址获得:float*ptr=&GlobalVar;CUDA编程框架/全局变量声明_host_,_device_._global_,

14、_constant_,_texture_/函数原型声明_global_ void kernelOne()/内核函数float handyFunction()/普通函数main()cudaMalloc(&d_GlblVarPtr,bytes)/在设备上分配空间cudaMemCpy(d_GlblVarPtr,h_Gl)/从主机传输数据到设备执行内核函数时的配置kernelOne(args);/调用内核函数cudaMemCpy(h_GlblVarPtr,);/从设备传输结果到主机_global_ void kernelOne(type args,)/内核函数/局部变量声明_local_,_shared_ /自动变量被默认分配到寄存器或本地存储器中float handyFunction(int inVar)/普通函数