搜索
GPUGPU简介GPU研究现状GPU内部架构CUDA编程模型CUDA存储器模型CUDA程序实例CUDA程序优化GPU简介单核处理器芯片已经到了尽头PowerWall功耗大MemoryWall存储器延迟很难降低，缓存占据70%芯片面积GPU简介多核和众核时代多个适当复杂度、低功耗核心并行工作时钟频率基本不变未来计算机硬件不会更快，但会更“宽”必须重新设计算法“Multicore:Thisistheonewhichwillhavethebiggestimpactonus.Wehaveneverhadaproblemtosolvelikethis.Abreakthroughisneededinhowapplicationsaredoneonmulticoredevices.”-BillGates,MicrosoftGPU简介GPU，GraphicsProcessingUnit的简写，是现代显卡中非常重要的一个部分，其地位与CPU在主板上的地位一致，主要负责的任务是加速图形处理速度。GPU是一个高度并行化的多线程、多核心处理器。GPU简介•GPU/CPU计算能力比较GPU简介•GPU/CPU存储器带宽比较GPU简介•市场迫切需要实时、高清晰度的3D图形，可编程的GPU已发展成为一种高度并行化、多线程、多核心的处理器，具有杰出的计算能力和极高的存储器带宽。•GPU极大提升了计算机图形处理的速度、增强了图形的质量，并促进了与计算机图形相关其他应用领域的快速发展。与中央处理器(CentralProcessingUnit，CPU)的串行设计模式不同，GPU为图形处理设计，具有天然的并行特性。GPU简介型号适用用户GeForce家庭和企业的娱乐应用，面向游戏用户Quadro应用于图形工作站，面向专业级用户Tesla用于高性能通用计算，面向研究人员Tegra适用于移动设备Ion上网本GPU研究现状•NVIDIA在1999年推出了第一款GPU产品－GeForce256。主要任务是进行图形渲染任务，缓解CPU压力。•从GPU诞生那天开始，其发展脚步就没有停止下来，由于其独特的体系架构和超强的浮点运算能力，人们希望将某些通用计算问题移植到GPU上来完成以提升效率，出现了所谓的GPGPU（GeneralPurposeGraphicProcessUnit），但是由于其开发难度较大，没有被广泛接受。•2006年NVIDIA推出了第一款基于Tesla架构的GPU（G80），GPU已经不仅仅局限于图形渲染，开始正式向通用计算领域迈进。GPU研究现状•2007年6月，NVIDIA推出了CUDA（ComputerUnifiedDeviceArchitecture计算统一设备结构）。CUDA是一种将GPU作为数据并行计算设备的软硬件体系。在CUDA的架构中，不再像过去GPGPU架构那样将通用计算映射到图形API中，对于开发者来说，CUDA的开发门槛大大降低了。CUDA的编程语言基于标准C，因此任何有C语言基础的用户都很容易地开发CUDA的应用程序。由于这些特性，CUDA在推出后迅速发展，被广泛应用于石油勘测、天文计算、流体力学模拟、分子动力学仿真、生物计算、图像处理、音视频编解码等领域。GPU内部架构CPU:强控制弱计算，更多资源用于缓存CPU:强计算弱控制，更多资源用于数据计算GPU内部架构•GPU体系架构在不断的发展，以GT200体系架构为代表对GPU的并行层次进行分析。TeslaGT200由两部分组成，分别是可伸缩流处理器阵列(ScalableStreamingProcessorArray，SPA)和存储器系统，它们由一个片上互联网络连接。如下图所示，可伸缩流处理器阵列由若干个线程处理器群(ThreadProcessingCluster，TPC)构成，每个TPC包含2～3个流多处理器(StreamingMultiprocessor，SM)，每个流多处理器中包含8个流处理器(StreamingProcessor，SP)。流处理器有独立的寄存器和指令指针，但缺少取指和调度单元，而流多处理器才拥有完整前端，包括取值、译码、发射等。从结构上看，每个流多处理器相当于一个8路单指令流多数据流(SingleInstructionMultipleData，SIMD)处理器，不同的是，GPU实现了自动向量机化，NVIDIA将之命名为单指令流多线程(SingleInstructionMultipleThread，SIMT)GPU内部架构前进GPU内部架构•在GPU中，流多处理器才能被称为真正的完整核心，整个可伸缩流处理器阵列可以被看成是由多个流多处理器组成的多单指令流多线程(MultipleSIMT，MSIMT)系统。TeslaGT200架构在可编程性和灵活性与硬件的复杂度和功耗之间取得了很好的折衷，线程被组织成多个线程块(ThreadBlock)，分配到各个流多处理器上，而每个线程块内的线程再被以单指令流多线程的方式交给流处理器运行。GPU内部架构•由于CPU和GPU设计目标的不同导致了两者在架构、并行层次和性能方面差异较大:CPU的重线程与GPU的轻线程CPU的MIMD多核与GPU的SIMT众核（x7560）CPU内存、缓存与GPU存储器•GPU是以大量线程实现面向吞吐量的数据并行计算，适合于处理计算密度高、逻辑分支简单的大规模数据并行负载；而CPU则有复杂的控制逻辑和大容量的缓存减小延迟，擅长复杂逻辑运算。GPU编程•CUDA•OpenCLOpenCL(全称OpenComputingLanguage，开放运算语言)是第一个面向异构系统通用目的并行编程的开放式、免费标准，也是一个统一的编程环境，便于软件开发人员为高性能计算服务器、桌面计算系统、手持设备编写高效轻便的代码，而且广泛适用于多核心处理器(CPU)、图形处理器(GPU)。CUDA编程模型•CUDA的基本思想是支持大量的线程级并行（ThreadLevelParallel），并在硬件中动态地调度和执行这些线程。•异构思想Host+Device(Coprocessor)让GPU来运行一些能够被高度线程化的代码。GPU与CPU协同工作，GPU只有在计算高度数据并行任务时才发挥作用。CUDA编程模型•调用kernel函数时CPU调用API将显卡端程序的二进制代码传到GPU•grid运行在SPA上•block运行在SM上•thread运行在SP上返回CUDA编程模型•gridblockthreadKernel不是一个完整的程序，而只是其中的一个关键并行计算步骤Kernel以一个网格(Grid)的形式执行，每个网格由若干个线程块（block）组成，每一个线程块又由最多512个线程(thread)组成。一个grid最多可以有65535*65535个block一个block总共最多可以有512个thread，在三个维度上的最大值分别为512,512和64CUDA编程模型grid之间通过globalmemory交换数据block之间不能相互通信,只能通过globalmemory共享数据,不要让多个block写同一区段内容（不保证数据一致性和顺序一致性）同一block内的thread可以通过sharedmemory和同步实现通信block间粗粒度并行，block内thread细粒度并行CUDA编程模型•warpwarp是硬件特性带来的概念，在CUDAC语言中是透明的（除vote函数），但应用中不能忽略一个warp中有32个线程，这是因为SM中有8个SP，执行一条指令的延迟是4个周期，使用了流水线技术一个halfwarp中有16个线程，这是因为执行单元的频率是其他单元的两倍，每两个周期才进行一次数据传输CUDA编程模型•分支性能与现代的微处理器不同，NVIDIA的SM没有预测执行机制-没有分支预测单元(BranchPredicator)。在需要分支时，只有当warp中所有的线程都计算出各自的分支的地址，并且完成取指以后，warp才能继续往下执行。如果一个warp内需要执行N个分支，那么SM就需要把每一个分支的指令发射到每一个SP上，再由SP根据线程的逻辑决定需不需要执行。这是一个串行过程，此时SIMT完成分支的时间是多个分支时间之和。CUDA存储器模型•Register•Local•Shared•Global•Constant•Texture•Hostmemory•PinnedhostmemoryCUDA存储器模型•CUDA对C的扩展：kernel执行参数运算符，用来传递一些kernel执行参数Grid的大小和维度Block的大小和维度外部声明的sharedmemory大小stream编号CUDA存储器模型•执行参数与内建变量的作用各个thread和block之间的唯一不同就是threadID和BlockID,通过内建变量控制各个线程处理的指令和数据CPU运行核函数时的执行参数确定GPU在SPA上分配多少个block，在SM上分配多少个threadCUDA存储器模型•CUDAdriverAPI•CUDAruntimeAPICUDA程序模板main(){//AllocatememoryonGPUfloat*Md;cudaMalloc((void**)&Md,size);//CopydatafromCPUtoGPUcudaMemcpy(Md,M,size,cudaMemcpyHostToDevice);//CallGPUkernelfunctionkerneldimGrid,dimBlock(arguments);//CopydatafromGPUbacktoCPUCopyFromDeviceMatrix(M,Md);//FreedevicematricesFreeDeviceMatrix(Md);}CUDA程序实例•矩阵乘法矩阵乘法时间复杂度是O(abc)，其中abc分别表示两个矩阵大小是a×b以及b×c.•实验环境CPU:Intel(R)Xeon(R)E54302.66GHz(8核)内存:DDR313334GGPU:TeslaC2050(448核)显存:GDDR53GCUDA程序实例•测试数据a=b=c=1024a=b=c=2048•算法CPU单线程CPU多线程GPUCUDA程序实例•实验结果1024*10242048*2048CPU单线程4568.300781ms557428.375000msCPU多线程606.676025ms193791.609375msGPU17.434999ms113.415001加速比351174CUDA程序优化activeblock每个SM最多可以有768（G8x,G9x）或者1024(GT200)个activethread这些activethread最多可以属于8个block还有受到SM中sharedmemory和register的制约最后的activeblock数量是由以上四个条件中的“短板”决定CUDA程序优化指令优化•选用计算复杂度较小的算法，处理字长为32bit•并行度高，粗粒度并行多，细粒度并行有局部性•分支映射成固定运算，展开代码避免循环•在精度允许的前提下使用带有__前缀的快速算法，只在必要的部分使用双精度和64-bitint•使用移位运算代替整数除法和求余•使用vote,atomic,red等intrinsic函数实现算法•只在线程间通信前进行同步CUDA程序优化存储器访问优化•对显存的I/O成本很高，提高两次对显存访问之间的计算量，通过同时计算与访存隐藏延时•把适合的数据放入纹理和常数缓存等缓解带宽压力，提高读取速度•避免bankconflict，非合并访问或cpu-gpu数据传输使用数组的结构体,而不是结构体数组•使用对齐，类型转换等手段实现合并访问