《CUDA编程模型(入门)》由会员分享，可在线阅读，更多相关《CUDA编程模型(入门)（25页珍藏版）》请在金锄头文库上搜索。
1、CUDA编程模型,华南理工大学 陈虎 博士 ,GPU与CPU的差异,GPU(Graphics Process Unit),面向计算密集型和大量数据并行化的计算 大量的晶体管用于计算单元,通用CPU,面向通用计算 大量的晶体管用于Cache和控制电路,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),
2、第二代GPU(GeForece 8800),GeForce 8800的主要技术参数,16个流多处理器(SM) 每个SM中包含了8个流处理器SP 每个SP包含一个乘加单元 每个SM管理了24个线程簇(warp)，共有768个线程 采用单线程多数据(Single-thread, Multiple-data)技术 每个周期在8个SP上并行执行一个线程簇,GF8800的存储层次结构,GeForce 8800,GPU最适合做什么？,对多个数据进行同一种运算（STMD适用） 一次存储器访问，多次运算（外部DDR访问开销高，局部存储器容量较小） 浮点计算比例高（特别是单精度浮点） 典型计算：物理模拟，线性代
3、数计算 应用领域： 计算生物学 图像处理,CUDA编程模型,Nvidia公司开发的编程模型 可以和VC 8.0集成使用 下载地址： http:/ CUDA Application,PTX 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 runt
4、ime library (cudart) The CUDA core library (cuda),CUDA中的线程,线程： CUDA中的基本执行单元； 硬件支持，开销很小； 所有线程执行相同的代码（STMD） 每个线程有一个唯一的标识IDthreadIdx 若干线程还可以组成块(Block) 线程块可以呈一维、二维或者三维结构,CUDA存储器分配,cudaMalloc() 在全局存储器中分配空间 两个参数： 地址指针 空间大小 cudaFree() 在全局存储器中回收空间 参数 回收空间地址指针,Grid,Global Memory,Block (0, 0),Shared Memory,Th
5、read (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(floa
6、t);cudaMalloc(void*),主机和设备之间的数据传输,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 单精
7、度浮点数组 M在主机存储器中 Md在设备存储器中 传输方向常数： cudaMemcpyHostToDevice cudaMemcpyDeviceToHost,cudaMemcpy(Md.elements, M.elements, size, cudaMemcpyHostToDevice);cudaMemcpy(M.elements, Md.elements, size, cudaMemcpyDeviceToHost);,CUDA的函数声明,_global_定义一个内核函数 必须返回 void_device_ 和 _host_可以一起使用,CUDA 的函数声明(续),_device_ 函数不能使
8、用函数指针； _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
9、 1.0调用的内核函数都是异步的, 阻塞需要显式的同步。,CUDA变量类型修饰符,当使用_local_, _shared_, 和 _constant_时， _device_ 是可选的没有任何修饰符的自动变量将默认放在寄存器内 数组除外，它们将会被摆放在本地存储器内,20,在哪里声明变量?,全局存储器 常存储器,寄存器 (自动) 共享存储器 本地存储器,21,G80 的CUDA存储器实现,每一个thread 可以: 可以读写寄存器 可以读写本地存储器 每一个block可以读写共享存储器 每一个grid 可以读写全局存储器 每一个grid 只可读常数存储器,Grid,Global Memory,B
10、lock (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) 它的访问速度比共享存储器要慢 在设备上执行运算的优化方法是： 把数据分割成子集，放入共享存储器中 以一个线程块操作一个数据子集：从全局存储器读入数据子集到共享存储器，使用多个线程实现存储级并行 在共享存
11、储器中对数据子集执行运算; 每一个线程可以有效率地多次访问任何数据元素 从共享存储器上拷贝结果到全局存储器,通用的编程策略(续),常数存储器也在设备内存里(DRAM) 比访问共享存储器慢得多 但是它可以高速缓存，对于只读数据的访问非常高效。 通过访问模式小心划分数据 只读 常数存储器(若在高速缓存内，非常快) 在块内读写共享 共享存储器(非常快) 线程可读写 寄存器 (非常快) 读写输入和结果 全局存储器(非常慢),24,变量类型约束,指针只能指向在全局存储器已分配或声明的内存 : 在主机分配，然后传递给内核: _global_ void KernelFunc(float* ptr) 以全局变
12、量的地址获得: float* ptr = ,CUDA编程框架,/全局变量声明 _host_ , _device_. _global_, _constant_, _texture_ /函数原型声明 _global_ void kernelOne() /内核函数 float handyFunction() /普通函数 main()cudaMalloc( /从设备传输结果到主机 _global_ void kernelOne(type args,) /内核函数/局部变量声明_local_, _shared_ /自动变量被默认分配到寄存器或本地存储器中 float handyFunction(int inVar) /普通函数 ,
