基于CUDA的GPU并行计算.pptx

基于CUDA的GPU并行计算.pptx

《基于CUDA的GPU并行计算.pptx》由会员分享，可在线阅读，更多相关《基于CUDA的GPU并行计算.pptx（38页珍藏版）》请在冰豆网上搜索。

基于CUDA的GPU并行计算,ComputeUnifiedDeviceArchitecture,*,南京大学多媒体研究所,2,内容,概述CUDA的硬件CUDA的软件性能提升并行计算,概述,摩尔定律,*,南京大学多媒体研究所,4,摩尔定律的失效,晶体管尺寸流水线级数能量墙存储墙,*,南京大学多媒体研究所,5,新摩尔定律,未来计算机硬件不会更快，但会更“宽”Nolongergetfaster,justwider必须重新设计算法,*,南京大学多媒体研究所,6,1Basedonslide7ofS.Green,“GPUPhysics,”SIGGRAPH2007GPGPUCourse.http:

/www.gpgpu.org/s2007/slides/15-GPGPU-physics.pdf,Whatisdrivingthemany-cores?

GPU与CPU的不同,TheGPUisspecializedforcompute-intensive,massivelydataparallelcomputation（exactlywhatgraphicsrenderingisabout）So,moretransistorscanbedevotedtodataprocessingratherthandatacachingandflowcontrol,Control,ALUALU,DRAMDRAMThefast-growingvideogameindustryexertsstrongeconomicpressureforconstantinnovation,CALUALUPCaUche,GPU,*,南京大学多媒体研究所,9,什么是GPU计算,GLSL/HLSL,*,南京大学多媒体研究所,10,什么是CUDA,编程语言CUDA可以使用C、Fortan等语言，SL是类C语言执行方式CUDA编译成汇编语言，执行时由显卡编译为机器指令；SL则是在执行时载入、编译，速度较慢。

计算精度CUDA满足IEEE-754浮点标准，支持单、双精度浮点运算编程方式CUDA更加自由；SL则需要模拟渲染过程。

对硬件的适应,*,南京大学多媒体研究所,11,为什么GPU比CPU快,更高浮点运算能力更快的存储访问速度更多的快速存储单元更迅速的线程切换速度,*,南京大学多媒体研究所,12,更快的存储访问速度,CPU与内存通过总线进行数据交换GPU与其内存都位于显卡上,CUDA的硬件,G80CUDAmodeADeviceExample,Load/store,GlobalMemory,ThreadExecutionManager,HostInputAssembler,Texture,Texture,Texture,Texture,Texture,Texture,Texture,T,T,ex,x,t,tu,u,r,r,e,e,Load/store,Load/store,Load/store,Load/store,Load/store,StreamingMultiprocessor（SM）,*,南京大学多媒体研究所,15,CUDA执行模型,t0t1t2tm,Blocks,SP,SP,MSheamroeryd,MSheamroedry,MTIU,MTIU,t0t1t2tmBlocks,SM0SM1,Upto8blockstoeachSMasresourceallowsSMinG80cantakeupto768threads,存储器模型,RegisterLocalsharedGlobalConstantTextureHostmemoryPinnedhostmemory,寄存器与localmemory,对每个线程来说，寄存器都是线程私有的-这与CPU中一样。

如果寄存器被消耗完，数据将被存储在本地存储器（localmemory）。

Localmemory对每个线程也是私有的，但是localmemory中的数据是被保存在显存中，而不是片内的寄存器或者缓存中，速度很慢。

线程的输入和中间输出变量将被保存在寄存器或者本地存储器中。

Sharedmemory,用于线程间通信的共享存储器。

共享存储器是一块可以被同一block中的所有thread访问的可读写存储器。

访问共享存储器几乎和访问寄存器一样快，是实现线程间通信的延迟最小的方法。

共享存储器可以实现许多不同的功能，如用于保存共用的计数器（例如计算循环次数）或者block的公用结果（例如计算512个数的平均值，并用于以后的计算）。

constantmemory,texturememory,利用GPU用于图形计算的专用单元发展而来的高速只读缓存速度与命中率有关，不命中时将进行对显存的访问常数存储器空间较小（只有64k），支持随机访问。

从host端只写，从device端只读纹理存储器尺寸则大得多，并且支持二维寻址。

（一个数据的“上下左右”的数据都能被读入缓存）适合实现图像处理算法和查找表,全局存储器,使用的是普通的显存，无缓存，可读写，速度慢整个网格中的任意线程都能读写全局存储器的任意位置，并且既可以从CPU访问，也可以从CPU访问。

各种存储器的延迟,register:

1周期sharedmemory:

1周期（无bankconflict）-16周期（发生16路bankconflict）texturememory:

1（命中）-数百周期（不命中）constantmemory:

1（命中）-数百周期（不命中）globallocalmemory:

数百周期,各存储器大小,每个SM中有64K（GT200）或者32K（G8x,G9x）寄存器，寄存器的最小单位是32bit的registerfile每个SM中有16Ksharedmemory一共可以声明64K的constantmemory，但每个SM的cache序列只有8K可以声明很大的texturememory，但是实际上的texturecache序列为每SM6-8K,CUDA的软件,当前的GPU开发环境,Cg：

优秀的图形学开发环境，但不适合GPU通用计算开发ATIstream：

硬件上已经有了基础，但只有低层次汇编能够使用所有资源。

高层次抽象Brook本质上是基于上一代GPU的，缺乏良好的编程模型OpenCL：

联合制定的标准，抽象层次较低，对硬件直接操作更多，代码需要根据不同硬件优化CUDA：

目前最佳选择,下载CUDA软件,http:

/htmlCUDASDKCUDAVisualProfiler,*,南京大学多媒体研究所,26,CUDA程序模板,Main（）/AllocatememoryonGPUfloat*Md;cudaMalloc（void*）,*,南京大学多媒体研究所,27,性能提升,共享内存的使用,*,南京大学多媒体研究所,29,优化原则：

activeblock,一个SM中可以有多个block等待处理，在一个warp需要访问存储器或者同步时，另外一个warp可以使用执行单元的资源增加activeblock对提高SM利用率有好处增加activeblock只是手段，不是最终的评价标准。

最终目的是要隐藏延迟,优化原则：

activeblock,每个SM最多可以有768（G8x,G9x）或者1024（GT200）个activethread这些activethread最多可以属于8个block还有受到SM中sharedmemory和register的制约最后的activeblock数量是由以上四个条件中的“短板”决定,并行计算,VectorReductionwithBranchDivergence,0,1,2,3,4,5,6,9,7,8,10,11,2+3,4+5,6+7,8+910+11,4.7,8.11,8.15,0+11,0.2.3,03.7,Thread0,Thread2,Thread4,Thread6,Thread8,Thread10,Asimpleimplementation,AssumewehavealreadyloadedarrayintosharedfloatpartialSum,unsignedintt=threadIdx.x;for（unsignedintstride=1;strideblockDim.x;stride*=2）syncthreads（）;if（t%（2*stride）=0）partialSumt+=partialSumt+stride;,NoDivergenceuntil16sub-sums,Thread0,0,1,2,3,13,14,17,15,16,18,19,15+31,0+116,3,4,Abetterimplementation,AssumewehavealreadyloadedarrayintosharedfloatpartialSum,unsignedintt=threadIdx.x;for（unsignedintstride=blockDim.x;stride1;stride1）syncthreads（）;if（tstride）partialSumt+=partialSumt+stride;,MemoryCoalescing,Whenaccessingglobalmemory,peakperformanceutilizationoccurswhenallthreadsinaWarpaccesscontinuousmemorylocations.,Md,Nd,W,I,D,T,H,WIDTH,Thread1Thread2,Nocoalesced,coalesced,MemoryAccessPattern,MdNd,WIDTH,WIDTH,Md,Nd,OriginalAccessPattern,TiledAccessPattern,Copyintoscratchpadmemory,Performmultiplicationwithscratchpad,values,