CUDA超大規(guī)模并行程序設計.ppt

CUDA超大規(guī)模并行程序設計,鄧仰東dengyd清華大學微電子學研究所,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA基礎CUDA存儲器CUDA程序設計工具程序優(yōu)化,GraphicProcessingUnit(GPU),用于個人計算機、工作站和游戲機的專用圖像顯示設備顯示卡nVidia和ATI(nowAMD)是主要制造商Intel準備通過Larrabee進入這一市場主板集成Intel,3維圖像流水線,Framebuffer,Texture,CPU,GPU,實時3維高速圖形處理,一幀典型圖像1Mtriangles3Mvertices25Mfragments30frames/s30Mtriangles/s90Mvertices/s750Mfragments/s,傳統(tǒng)GPU架構,Graphicsprogram,Vertexprocessors,Fragmentprocessors,Pixeloperations,Outputimage,GPU的強大運算能力,數(shù)據(jù)級并行:計算一致性,專用存儲器通道有效隱藏存儲器延時,GeneralPurposeComputingonGPU(GPGPU),GPGPU,核心思想用圖形語言描述通用計算問題把數(shù)據(jù)映射到vertex或者fragment處理器但是硬件資源使用不充分存儲器訪問方式嚴重受限難以調試和查錯高度圖形處理和編程技巧,NVidiaG200Architecture,CUDA:ComputeUnifiedDeviceArchitecture,通用并行計算模型單指令、多數(shù)據(jù)執(zhí)行模式(SIMD)所有線程執(zhí)行同一段代碼(1000sthreadsonthefly)大量并行計算資源處理不同數(shù)據(jù)隱藏存儲器延時提升計算通信比例合并相鄰地址的內存訪問快速線程切換1cycleGPUvs.1000cyclesCPU,混合計算模型,CUDA:集成CPU+GPUC應用程序CPU:順序執(zhí)行代碼GPU=超大規(guī)模數(shù)據(jù)并行協(xié)處理器“批發(fā)”式執(zhí)行大量細粒度線程,kernel0,CPUSerialCode,CPUSerialCode,GPUParallelCode,GPUParallelCode,Concurrentexecution!,kernel1,CUDA成功案例,CUDA性能,BLAS3:127GFLOPS/基本線性代數(shù):matrix-matrixFFT:52benchFFT*GFLOPSFDTD:1.2Gcells/sec/計算電動力學SSEARCH:5.2Gcells/sec/Smith-Waterman基因序列比較BlackScholes:4.7GOptions/sec/期權定價模型VMD:290GFLOPS/分子動力學圖形顯示,ProblemInstancesforSparseMatrixVectorProduct(SMVP),SPMVThroughputonGTX280,SMVPApplication:StaticTimingAnalysis,AdaptedfromRamalingam,A.et.al.AnAccurateSparseMatrixBasedFrameworkforStatisticalStaticTimingAnalysis.ICCAD.2006.,StaticTimingAnalysisResultsonGTX280,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA基礎CUDA存儲器CUDA程序設計工具程序優(yōu)化,并行性的維度,1維y=a+b/y,a,bvectors2維P=MN/P,M,Nmatrices3維CTorMRIimaging,=,并行線程組織結構,Thread:并行的基本單位Threadblock:互相合作的線程組CooperativeThreadArray(CTA)允許彼此同步通過快速共享內存交換數(shù)據(jù)以1維、2維或3維組織最多包含512個線程Grid:一組threadblock以1維或2維組織共享全局內存Kernel:在GPU上執(zhí)行的核心程序Onekernelonegrid,ParallelProgramOrganizationinCUDA,Thread,Threadblock,Grid,SP,Software,Hardware,SM,GPU,并行線程執(zhí)行,調用kernelfunction需要指定執(zhí)行配置Threads和blocks具有IDsthreadIdx:1D,2D,or3DblockIdx:1D,or2D由此決定相應處理數(shù)據(jù),_global_voidkernel(.);dim3DimGrid(3,2);/6threadblocksdim3DimBlock(16,16);/256threadsperblockkernel(.);,實例1:Element-WiseAddition,/CPUprogram/sumoftwovectorsaandbvoidadd_cpu(float*a,float*b,intN)for(intidx=0;idx(a,b,N);,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA基礎CUDA存儲器CUDA程序設計工具程序優(yōu)化,CUDAProcessingFlow,并行線程執(zhí)行,SM內以(warp即32threads)為單位并行執(zhí)行Warp內線程執(zhí)行同一條指令Half-warp是存儲操作的基本單位,Warp,GPU負載分配,Globalblockscheduler管理threadblock級并行從CPU獲得線程組織信息根據(jù)硬件結構分配threadblock到SM,StreamingMultiprocessor(SM),StreamingMultiprocessor執(zhí)行ThreadBlocks,線程以block為單位分配到SM視資源需求,一個SM分配至多8個blockSMinG80可以接受768個線程256(threads/block)*3blocks或128(threads/block)*6blocks,etc.線程并發(fā)(concurrently)運行SM分配并維護線程IDSM管理并調度線程,ThreadLifeCycle,Grid在GPU上啟動Threadblocks順序分配到SMs一般SM應有1threadblockSM把線程組織為warpsSM調度并執(zhí)行就緒的warpWarps和threadblocks執(zhí)行結束后釋放資源GPU繼續(xù)分發(fā)threadblocks,ExampleofHidingMemoryLatency,G80:執(zhí)行warp全部線程的一條指令需要8個時鐘cycle假定1globalmemoryaccess/8instructionsA400-cycleglobalmemorylatencyHowmanywarpsareneededtotoleratethelatency?,400cycles*1MEM/8cycles=50cyclesperinstructiononaverage50cycles/4cyclesperwarp=12.513warpstokeepanSMbusy,ArithmeticInstructionThroughput,4clockcyclesSingle-precisionfloating-pointadd,multiply,andmultiply-add,Integeradd,24-bitintegermultiplicationBitwiseoperations,compare,min,max,typeconversioninstruction;,Note.1.Awarpisissuedin4cycles2.Arithmeticopsarepipelined.3.Stillpossibletohave8opsreadyineachcycle.,ArithmeticInstructionThroughput,16clockcyclesReciprocal,reciprocalsquareroot,32-bitIntegermultiplicationOtherfunctionsarecombinationsoftheabovey/x=rcp(x)*y/20cyclesperwarpsqrt(x)=rcp(rsqrt(x)/32cyclesperwarpIntegerdivisionandmodulooperationarecostly!,浮點數(shù)精度,GT200之前的GPUIEEE-754FloatingPointStandard單精度浮點數(shù)GT200增加雙精度浮點數(shù)支持SP仍然只支持單精度浮點數(shù)每個SM配置一個雙精度浮點單元雙精度比單精度運算慢8-12倍,控制流(ControlFlow),同一warp內的分支語句可能執(zhí)行不同的指令路徑不同指令路徑的線程只能順序執(zhí)行每次執(zhí)行warp中一條可能的路徑N條指令路徑1/Nthroughput只需要考慮同一warp即可，不同warp的不同的指令路徑不具相關性G80上使用指令預測技術加速指令執(zhí)行,控制流(ControlFlow),常見情況:分支條件是threadID的函數(shù)時,容易導致divergenceExamplewithdivergence:If(threadIdx.x2)在threadblock產(chǎn)生兩條不同指令路徑Branchgranularity2)也在threadblock產(chǎn)生兩條不同指令路徑Branchgranularityisawholemultipleofwarpsize同一warp的所有線程具備相同指令路徑,線程同步,void_syncthreads();Barriersynchronization同步threadblock之內的所有線程避免訪問共享內存時發(fā)生RAW/WAR/WAW冒險(hazard),_shared_floatscratch256;scratchthreadID=beginthreadID;_syncthreads();intleft=scratchthreadID-1;,在此等待，直至所有線程到達才開始執(zhí)行下面的代碼,Dead-Lockwith_syncthreads,Dead-lockifSomethreadshavevallargerthanthresholdAndothersnot,_global_voidcompute(.)/dosomecomputationforvalif(valthreshold)return;_syncthreads();/workwithval,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA基礎CUDA存儲器CUDA程序設計工具程序優(yōu)化,CUDA擴展語言結構,Declspecsglobal,device,shared,local,constantKeywordsthreadIdx,blockIdxthreadDim,blockDimIntrinsics_syncthreadsRuntimeAPIMemory,symbol,executionmanagementFunctionlaunch,_device_floatfilterN;_global_voidconvolve(float*image)_shared_floatregionM;.regionthreadIdx=imagei;_syncthreads().imagej=result;/AllocateGPUmemoryvoid*myimage=cudaMalloc(bytes)/100blocks,10threadsperblockfoo(parameters);,存儲器空間,R/Wper-threadregisters1-cyclelatencyR/Wper-threadlocalmemorySlowregisterspillingtoglobalmemoryR/Wper-blocksharedmemory1-cyclelatencyButbankconflictsmaydragdownR/Wper-gridglobalmemory500-cyclelatencyButcoalescingaccessingcouldhidelatencyReadonlyper-gridconstantandtexturememories500-cyclelatencyButcached,GPUGlobalMemory分配,cudaMalloc()分配顯存中的globalmemory兩個參數(shù)對象數(shù)組指針和數(shù)組尺寸cudaFree()釋放顯存中的globalmemory對象數(shù)組指針,intblk_sz=64;float*Md;intsize=blk_sz*blk_sz*sizeof(float);cudaMalloc(void*),HostDevice數(shù)據(jù)交換,cudaMemcpy()MemorydatatransferRequiresfourparametersPointertodestinationPointertosourceNumberofbytescopiedTypeoftransferHosttoHost,HosttoDevice,DevicetoHost,DevicetoDevice,cudaMemcpy(Md,M.elements,size,cudaMemcpyHostToDevice);cudaMemcpy(M.elements,Md,size,cudaMemcpyDeviceToHost);,CUDA引入的新變量類型,_device_儲存于GPU上的globalmemory空間和應用程序具有相同的生命期(lifetime)可被grid中所有線程存取,CPU代碼通過runtime函數(shù)存取_constant_儲存于GPU上的constantmemory空間和應用程序具有相同的生命期(lifetime)可被grid中所有線程存取,CPU代碼通過runtime函數(shù)存取_shared_儲存于GPU上threadblock內的共享存儲器和threadblock具有相同的生命期(lifetime)只能被threadblock內的線程存取Local變量儲存于SM內的寄存器和localmemory和thread具有相同的生命期(lifetime)Thread私有,CUDA函數(shù)定義,_global_定義kernel函數(shù)必須返回void_device_函數(shù)不能用intheight;float*elements;Matrix;,A,B,C,WM.width=N.heightI,M.height,M.width,N.width,實例2:矩陣相乘,C=ABofsizeWIDTHxWIDTH一個線程處理一個矩陣元素簡化:假定WIDTHxWIDTH(Ad,Bd,Cd,M.width);/ReadPfromthedevicecopyFromDeviceMatrix(C.elements,Cd);cudaMemCopy(C,Cd,N*size,cudaMemcpyDeviceToHost);/FreedevicematricescudaFree(Ad);cudaFree(Bd);cudaFree(Cd);,CUDAImplementationKernel,/Matrixmultiplicationkernelthreadspecification_global_voidMuld(float*Ad,float*Bd,float*Cd,intwidth)/2DThreadIDinttx=threadIdx.x;intty=threadIdx.y;/cvalueisusedtostoretheelementofthematrix/thatiscomputedbythethreadfloatcvalue=0;,CUDAImplementationKernel,A,B,C,WIDTH,WIDTH,WIDTH,WIDTH,ty,tx,for(intk=0;kwidth;+k)floatae=Adty*width+k;floatbe=Bdtx+k*width;cvalue+=ae*be;/Writethematrixtodevicememory;/eachthreadwritesoneelementCdty*width+tx=cvalue;,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA存儲器SharedmemoryGlobalmemoryCUDA程序設計工具程序優(yōu)化,共享存儲器(SharedMemory),設置于streamingmultiprocessor內部由一個線程塊內部全部線程共享完全由軟件控制訪問一個地址只需要1個時鐘周期,共享存儲器結構,G80的共享存儲器組織為16banksAddressedin4bytesBankID=4-byteaddress%16相鄰4-byte地址映射相鄰banks每一bank的帶寬為4bytesperclockcycle對同一bank的同時訪問導致bankconflict只能順序處理僅限于同一線程塊內的線程,BankAddressing實例,NoBankConflictsLinearaddressingstride=1(s=1),NoBankConflictsRandom1:1Permutation,_shared_floatshared256;floatfoo=sharedthreadIdx.x;,BankAddressing實例,2-waybankconflictsLinearaddressingstride=2(s=2),8-waybankconflictsLinearaddressingstride=8(s=8),_shared_floatshared256;floatfoo=shared2*threadIdx.x;,_shared_floatshared256;floatfoo=shared8*threadIdx.x;,常見BankConflict模式,Sharedmemory存放2D浮點數(shù)組16x16-elelmentsharedmemory1個線程處理矩陣的一行循環(huán)處理一行16個元素同一block的線程同時訪問一列即column1inpurple16-waybankconflicts,BankIndiceswithoutPadding,Bank,t15,解決方案,方案1:padtherows在每行最后添加一個元素方案2:transposebeforeprocessingSufferbankconflictsduringtransposeButpossiblysavethemlater,BankIndiceswithPadding,Transpose,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA存儲器SharedmemoryGlobalmemoryCUDA程序設計工具程序優(yōu)化,全局內存(GlobalMemory),全局內存在G80上沒有緩存Constantmemory和texturememory有少量緩存存取延時400-600clockcycles非常容易成為性能瓶頸優(yōu)化是提高性能的關鍵!,CoalescedGlobalMemoryAccesses,在half-warp層次對訪問globalmemory進行協(xié)調訪問連續(xù)globalmemory區(qū)域:64bytes-eachthreadreadsaword:int,float,128bytes-eachthreadreadsadouble-word:int2,float2,256byteseachthreadreadsaquad-word:int4,float4,額外限制:Globalmemory區(qū)域的起始地址必須是該區(qū)域數(shù)據(jù)類型尺寸的整數(shù)倍Warp中第k個線程訪問第k個地址例外:可以有某些中間線程不參加Predicatedaccess,divergencewithinawarp,CoalescedGlobalMemoryAccesses,Non-CoalescedGlobalMemoryAccesses,Non-CoalescedGlobalMemoryAccesses,提綱,從GPGPU到CUDACUDA并行程序組織并行執(zhí)行模型CUDA存儲器SharedmemoryGlobalmemoryCUDA程序設計工具程序優(yōu)化,下載CUDA軟件,軟件環(huán)境,GPU硬件和CUDA軟件安裝后:,CPU(Host),CUDALibraries(CUFFTfloat3a=d_inindex;a.x+=2;a.y+=2;a.z+=2;d_outindex=a;,Uncoalescedfloat3Code,float3需要12bytes:float3f=d_inthreadIdx.x;Eachthreadendsupexecuting3readssizeof(float3)4,8,or16Half-warpreadsthree64Bnon-contiguousregions,Coalescingfloat3Access,A3-stepapproach(256threads/block),GlobalMemory,Sharedmemory,Sharedmemory,Coalescingfloat3Access,Usesharedmemorytoallowcoalescing256threadsperblockAthreadblockneedssizeof(float3)x256bytesofSMEMEachthreadreads3scalarfloats:Offsets:0,(threads/block),2*(threads/block)Thesewilllikelybeprocessedbyotherthreads,sosyncProcessingEachthreadretrievesitsfloat3fromSMEMarrayCasttheSMEMpointerto(float3*)UsethreadIDasindexRestofthecomputecodedoesnotchange!,Coalescingfloat3Access代碼,MatrixTranspose,SDKSample(“transpose”)解釋通過sharedmemory實現(xiàn)coalescing在小尺度數(shù)據(jù)即可顯示優(yōu)化的明顯效果,UncoalescedTranspose,_global_voidtranspose_naive(float*odata,float*idata,intwidth,intheight)unsignedintxIndex=blockDim.x*blockIdx.x+threadIdx.x;unsignedintyIndex=blockDim.y*blockIdx.y+threadIdx.y;if(xIndexwidth,tx,ty,Height,Width,Height,Width,UncoalescedTranspose,CoalescedTranspose,假設:矩陣已被分解為方塊(tile)Threadblock(bx,by):Readthe(bx,by)inputtile,storeintoSMEMWritetheSMEMdatato(by,bx)outputtileThread(tx,ty):Readselement(tx,ty)frominputtileWriteselement(tx,ty)intooutputtileCoalescingisachievedif:Block/tiledimensionsaremultiplesof16,CoalescedTranspose,CoalescedTranspose,_global_voidtranspose(float*odata,float*idata,intwidth,intheight)_shared_floatblockBLOCK_DIM*BLOCK_DIM;unsignedintxBlock=blockDim.x*blockIdx.x;unsignedintyBlock=blockDim.y*blockIdx.y;unsignedintxIndex=xBlock+threadIdx.x;