CUDA编程模型(入门)

单击此处编辑母版标题样式,单击此处编辑母版文本样式,第二级,第三级,第四级,第五级,*,CUDA,编程模型,华南理工大学 陈虎 博士,GPU,与,CPU,的差异,GPU(Graphics Process Unit),面向计算密集型和大量数据并行化的计算,大量的晶体管用于计算单元,通用,CPU,面向通用计算,大量的晶体管用于,Cache,和控制电路,DRAM,Cache,ALU,Control,ALU,ALU,ALU,DRAM,CPU,GPU,GPU,与,CPU,的峰值速度比较,1,Based on slide 7 of S.Green,“GPU Physics,”SIGGRAPH 2007 GPGPU Course.http:/www.gpgpu.org/s2007/slides/15-GPGPU-physics.pdf,第一代,GPU,结构,(GeForce 6800),第二代,GPU(GeForece 8800),GeForce 8800,的主要技术参数,晶体管数目（百万）,681,工艺,90nm,芯片面积,(mm,2,),470,工作主频（,GHZ,）,1.5,峰值速度（,Gflops,）,576,处理器数目,128,片上存储器容量（,KB,）,488,功耗（,W,）,150,16,个流多处理器,(,SM,),每,个,SM,中包,含,了,8,个,流处理器,SP,每个,SP,包含一个乘加单元,每个,SM,管理了,24,个线程簇,(warp),，共有,768,个线程,采用单线程多数据,(Single-thread,Multiple-data),技术,每个周期在,8,个,SP,上并行执行一个线程簇,GF8800,的存储层次结构,层次,容量,(KB),延迟,(ns),局部存储器,16,26,L1 Cache,5,280,L2 Cache,32,370,DDR,510,GeForce 8800,GPU,最适合做什么？,对多个数据进行同一种运算（,STMD,适用）,一次存储器访问，多次运算（外部,DDR,访问开销高，局部存储器容量较小）,浮点计算比例高（特别是单精度浮点）,典型计算：物理模拟，线性代数计算,应用领域：,计算生物学,图像处理,CUDA,编程模型,Nvidia,公司开发的编程模型,可以和,VC 8.0,集成使用,下载地址：,http:/ to Target,Compiler,G80,GPU,Target code,PTX Code,Virtual,Physical,CPU Code,float4 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,）,每个线程有一个唯一的标识,ID,threadIdx,若干线程还可以组成块,(Block),线程块可以呈一维、二维或者三维结构,Device,Grid 1,Block,(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(),在全局存储器中回收空间,参数,回收空间地址指针,Grid,Global,Memory,Block(0,0),Shared Memory,Thread(0,0),Registers,Thread(1,0),Registers,Block(1,0),Shared Memory,Thread(0,0),Registers,Thread(1,0),Registers,Host,示范代码,分配,64*64,的单精度浮点数组,存储器地址为,Md.elements,TILE_WIDTH=64;,Matrix Md,/*“d”is often used to indicate a device data structure*/,int size=TILE_WIDTH*TILE_WIDTH*sizeof(float);,cudaMalloc(void*),cudaFree(Md.elements);,主机和设备之间的数据传输,cudaMemcpy(),存储器数据传输,参数：,目的地址,源地址,传输字节数,传输类型,主机,主机,主机,全局存储器,全局存储器,全局存储器,异步传输,Grid,Global,Memory,Block(0,0),Shared Memory,Thread(0,0),Registers,Thread(1,0),Registers,Block(1,0),Shared Memory,Thread(0,0),Registers,Thread(1,0),Registers,Host,示范代码,传输,64*64,单精度浮点数组,M,在主机存储器中,Md,在设备存储器中,传输方向常数：,cudaMemcpyHostToDevice,cudaMemcpyDeviceToHost,cudaMemcpy(Md.elements,M.elements,size,cudaMemcpyHostToDevice);,cudaMemcpy(M.elements,Md.elements,size,cudaMemcpyDeviceToHost);,CUDA,的函数声明,在哪里执行,:,只能从哪里调用,:,_device_,float DeviceFunc(),设备,设备,_global_,void KernelFunc(),设备,主机,_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,size_t SharedMemBytes=64;,/64 bytes of shared memory,KernelFunc,(.);,/,用,指明该函数调用的线程数,任何从,CUDA 1.0,调用的内核函数都是异步的,阻塞需要显式的同步。,CUDA,变量类型修饰符,当使用,_local_,_shared_,和,_constant_,时，,_device_,是可选的,没有任何修饰符的自动变量将默认放在寄存器内,数组除外，它们将会被摆放在本地存储器内,变量声明,存储器,作用域,生命期,_device_ _local_,int LocalVar;,本地存储器,thread,thread,_device_ _shared_,int SharedVar;,共享存储器,block,block,_device_,int GlobalVar;,全局存储器,grid,application,_device_ _constant_,int ConstantVar;,常存储器,grid,application,20,在哪里声明变量,?,主机能否访问,?,在任何函数外,在内核,可以,不可以,全局存储器,常存储器,寄存器,(,自动,),共享存储器,本地存储器,21,G80,的,CUDA,存储器实现,每一个,thread,可以,:,可以读写寄存器,可以读写本地存储器,每一个,block,可以读写,共享存储器,每一个,grid,可以读写,全局存储器,每一个,grid,只可读,常数存储器,Grid,Global Memory,Block(0,0),Shared Memory,Thread(0,0),Registers,Thread(1,0),Registers,Block(1,0),Shared Memory,Thread(0,0),Registers,Thread(1,0),Registers,Host,Constant Memory,22,通用的编程策略,全局存储器在设备内存中,(DRAM),它的访问速度比共享存储器要慢,在设备上执行运算的优化方法是：,把数据分割成子集，放入共享存储器中,以一个线程块操作一个数据子集：,从全局存储器读入数据子集到共享存储器，使用多个线程实现存储级并行,在共享存储器中对数据子集执行运算,;,每一个线程可以有效率地多次访问任何数据元素,从共享存储器上拷贝结果到全局存储器,通用的编程策略,(,续,),常数存储器也在设备内存里,(DRAM),比访问共享存储器慢得多,但是它可以高速缓存，对于只读数据的访问非常高效。,通过访问模式小心划分数据,只读 常数存储器,(,若在高速缓存内，非常快,),在块内读写共享 共享存储器,(,非常快,),线程可读写 寄存器,(,非常快,),读写输入和结果 全局存储器,(,非常慢,),24,变量类型约束,指针只能指向在全局存储器已分配或声明的内存,:,在主机分配，然后传递给内核,:,_global_ void KernelFunc(float*ptr),以全局变量的地址获得,:,float*ptr=,CUDA,编程框架,/,全局变量声明,_host_,_device_._global_,_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),/,普通函数,