GPUGPU简介GPU研究现状GPU内部架构CUDA编程模型CUDA存储器模型CUDA程序实例CUDA程序优化GPU简介单
核处理器芯片已经到了尽头PowerWall 功耗大MemoryWall 存储器延迟很难降低，缓存占据>70%芯片面积
GPU简介多核和众核时代多个适当复杂度、低功耗核心并行工作时钟频率基本不变未来计算机硬件不会更快，但会更“宽”必须重
新设计算法“Multicore:Thisistheonewhichwillhavethebiggestimp
actonus.Wehaveneverhadaproblemtosolvelikethis.Abre
akthroughisneededinhowapplicationsaredoneonmulticorede
vices.” -BillGates,MicrosoftGPU简介GPU，GraphicsPro
cessingUnit的简写，是现代显卡中非常重要的一个部分，其地位与CPU在主板上的地位一致，主要负责的任务是加速图形处
理速度。GPU是一个高度并行化的多线程、多核心处理器。GPU简介GPU/CPU计算能力比较GPU简介GPU/CPU存储
器带宽比较GPU简介市场迫切需要实时、高清晰度的3D图形，可编程的GPU已发展成为一种高度并行化、多线程、多核心的处理
器，具有杰出的计算能力和极高的存储器带宽。GPU极大提升了计算机图形处理的速度、增强了图形的质量，并促进了与计算机图形相关其他应
用领域的快速发展。与中央处理器(CentralProcessingUnit，CPU)的串行设计模式不同，GPU为图形处理设计，
具有天然的并行特性。GPU简介GPU研究现状NVIDIA在1999年推出了第一款GPU产品－GeForce256。主要任务
是进行图形渲染任务，缓解CPU压力。从GPU诞生那天开始，其发展脚步就没有停止下来，由于其独特的体系架构和超强的浮点运算能力，人
们希望将某些通用计算问题移植到GPU上来完成以提升效率，出现了所谓的GPGPU（GeneralPurposeGraphicP
rocessUnit），但是由于其开发难度较大，没有被广泛接受。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个流处理器(StreamingPr
ocessor，SP)。流处理器有独立的寄存器和指令指针，但缺少取指和调度单元，而流多处理器才拥有完整前端，包括取值、译码、发射等
。从结构上看，每个流多处理器相当于一个8路单指令流多数据流(SingleInstructionMultipleData，SI
MD)处理器，不同的是，GPU实现了自动向量机化，NVIDIA将之命名为单指令流多线程(SingleInstructionMu
ltipleThread，SIMT)GPU内部架构GPU内部架构在GPU中，流多处理器才能被称为真正的完整核心，整个可伸缩
流处理器阵列可以被看成是由多个流多处理器组成的多单指令流多线程(MultipleSIMT，MSIMT)系统。TeslaGT20
0架构在可编程性和灵活性与硬件的复杂度和功耗之间取得了很好的折衷，线程被组织成多个线程块(ThreadBlock)，分配到各个流
多处理器上，而每个线程块内的线程再被以单指令流多线程的方式交给流处理器运行。GPU内部架构由于CPU和GPU设计目标的不同导致
了两者在架构、并行层次和性能方面差异较大:CPU的重线程与GPU的轻线程CPU的MIMD多核与GPU的SIMT众核（x7560
）CPU内存、缓存与GPU存储器GPU是以大量线程实现面向吞吐量的数据并行计算，适合于处理计算密度高、逻辑分支简单的大规模数
据并行负载；而CPU则有复杂的控制逻辑和大容量的缓存减小延迟，擅长复杂逻辑运算。GPU编程CUDAOpenCL OpenC
L(全称OpenComputingLanguage，开放运算语言)是第一个面向异构系统通用目的并行编程的开放式、免费标准，也是
一个统一的编程环境，便于软件开发人员为高性能计算服务器、桌面计算系统、手持设备编写高效轻便的代码，而且广泛适用于多核心处理器(CP
U)、图形处理器(GPU)。CUDA编程模型CUDA的基本思想是支持大量的线程级并行（ThreadLevelParalle
l），并在硬件中动态地调度和执行这些线程。异构思想 Host+Device(Coprocessor) 让GPU来运行一些
能够被高度线程化的代码。GPU与CPU协同工作，GPU只有在计算高度数据并行任务时才发挥作用。CUDA编程模型调用kernel
函数时CPU调用API将显卡端程序的二进制代码传到GPUgrid运行在SPA上block运行在SM上thread运行在SP上
CUDA编程模型gridblockthread Kernel不是一个完整的程序，而只是其中的一个关键并行计算步骤 Ke
rnel以一个网格(Grid)的形式执行，每个网格由若干个线程块（block）组成，每一个线程块又由最多512个线程(thread
)组成。 一个grid最多可以有6553565535个block 一个block总共最多可以有512个thread，在三
个维度上的最大值分别为512,512和64CUDA编程模型 grid之间通过globalmemory交换数据 block
之间不能相互通信,只能通过globalmemory共享数据,不要让多个block写同一区段内容（不保证数据一致性和顺序一致性）
同一block内的thread可以通过sharedmemory和同步实现通信block间粗粒度并行，block内thread细
粒度并行CUDA编程模型warp warp是硬件特性带来的概念，在CUDAC语言中是透明的（除vote函数），但应用中不能
忽略 一个warp中有32个线程，这是因为SM中有8个SP，执行一条指令的延迟是4个周期，使用了流水线技术 一个halfwa
rp中有16个线程，这是因为执行单元的频率是其他单元的两倍，每两个周期才进行一次数据传输CUDA编程模型分支性能 与现代的
微处理器不同，NVIDIA的SM没有预测执行机制-没有分支预测单元(BranchPredicator)。 在需要分支时，只有
当warp中所有的线程都计算出各自的分支的地址，并且完成取指以后，warp才能继续往下执行。 如果一个warp内需要执行N个分
支，那么SM就需要把每一个分支的指令发射到每一个SP上，再由SP根据线程的逻辑决定需不需要执行。这是一个串行过程，此时SIMT完成
分支的时间是多个分支时间之和。CUDA存储器模型RegisterLocalSharedGlobalConstant
TextureHostmemoryPinnedhostmemoryCUDA存储器模型CUDA对C的扩展：kernel
执行参数 <<<>>>运算符，用来传递一些kernel执行参数 Grid的大小和维度 Block的大小和维度 外部声明
的sharedmemory大小 stream编号CUDA存储器模型执行参数与内建变量的作用 各个thread和block
之间的唯一不同就是threadID和BlockID,通过内建变量控制各个线程处理的指令和数据 CPU运行核函数时的执行参数确定G
PU在SPA上分配多少个block，在SM上分配多少个threadCUDA存储器模型CUDAdriverAPICUDA
runtimeAPICUDA程序模板main(){ //AllocatememoryonGPU floatM
d; cudaMalloc((void)&Md,size); //CopydatafromCPUtoGPU
cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice); //CallGPUkern
elfunction kernel<<>>(arguments); //Copyd
atafromGPUbacktoCPU CopyFromDeviceMatrix(M,Md); //Freede
vicematrices FreeDeviceMatrix(Md);}CUDA程序实例矩阵乘法矩阵乘法时间复杂度是O
(abc)，其中abc分别表示两个矩阵大小是a×b以及b×c.实验环境 CPU:Intel(R)Xeon(R)
E54302.66GHz(8核) 内存:DDR313334G GPU:TeslaC2050(448核) 显
存:GDDR53GCUDA程序实例测试数据a=b=c=1024a=b=c=2048算法CPU单线程CPU多线程
GPUCUDA程序实例实验结果CUDA程序优化activeblock 每个SM最多可以有768（G8x,G9x）或者1
024(GT200)个activethread 这些activethread最多可以属于8个block 还有受到SM中sh
aredmemory和register的制约 最后的activeblock数量是由以上四个条件中的“短板”决定CUDA程序
优化指令优化选用计算复杂度较小的算法，处理字长为32bit并行度高，粗粒度并行多，细粒度并行有局部性分支映射成固定运算，展
开代码避免循环在精度允许的前提下使用带有__前缀的快速算法，只在必要的部分使用双精度和64-bitint使用移位运算代替整数
除法和求余使用vote,atomic,red等intrinsic函数实现算法只在线程间通信前进行同步CUDA程序优化存储器
访问优化对显存的I/O成本很高，提高两次对显存访问之间的计算量，通过同时计算与访存隐藏延时把适合的数据放入纹理和常数缓存等缓解
带宽压力，提高读取速度避免bankconflict，非合并访问或cpu-gpu数据传输使用数组的结构体,而不是结构体数组使
用对齐，类型转换等手段实现合并访问1024102420482048CPU单线程4568.300781ms5574
28.375000ms加速比351174CPU多线程606.676025ms193791.609375msGPU1
7.434999ms113.415001英特尔（Intel）创始人之一戈登·摩尔，集成电路上可容纳的晶体管数目，约每隔18
个月便会增加一倍，性能也将提升一倍(为了不断提升处理速度，从处理器诞生的那天起到现在发展到今天已经经历了好几代的技术革新，比如
说对制作工艺，缓存，流水线技术等改进，但是当单核处理器发展到一定阶段，人们发现单纯的依靠增加处理器主频的方式提高处理器性能遇到
了瓶颈)，（一味的以增加单核处理器主频，散热量将异常大，甚至超过太阳表面的温度）（完成用户的任务，耗时图形）型号适用用户GeForce家庭和企业的娱乐应用，面向游戏用户Quadro应用于图形工作站，面向专业级用户Ion上网本Tegra适用于移动设备Tesla用于高性能通用计算，面向研究人员前进返回英特尔（Intel）创始人之一戈登·摩尔，集成电路上可容纳的晶体管数目，约每隔18个月便会增加一倍，性能也将提升一倍(为了不断提升处理速度，从处理器诞生的那天起到现在发展到今天已经经历了好几代的技术革新，比如说对制作工艺，缓存，流水线技术等改进，但是当单核处理器发展到一定阶段，人们发现单纯的依靠增加处理器主频的方式提高处理器性能遇到了瓶颈)，（一味的以增加单核处理器主频，散热量将异常大，甚至超过太阳表面的温度）（完成用户的任务，耗时图形） |
|
