cuda程序设计.ppt
- 文档编号:1368326
- 上传时间:2022-10-21
- 格式:PPT
- 页数:76
- 大小:2.97MB
cuda程序设计.ppt
《cuda程序设计.ppt》由会员分享,可在线阅读,更多相关《cuda程序设计.ppt(76页珍藏版)》请在冰豆网上搜索。
CUDA程序设计,主要内容,GPGPU及CUDA介绍CUDA编程模型多线程及存储器硬件,GPGPU及CUDA介绍,多核时代,多个适当复杂度、相对低功耗内核并行工作配置并行硬件资源提高处理能力核心时钟频率基本不变,nVidiaGT200,Quad-coreOpteron,IBMCellBroadbandEngine,GPU与CPU硬件架构的对比,CPU:
更多资源用于缓存及流控制GPU:
更多资源用于数据计算适合具备可预测、针对数组的计算模式,CPU,GPU,应用范围,CPU:
controlprocessor不规则数据结构不可预测存取模式递归算法分支密集型算法单线程程序,GPU:
dataprocessor规则数据结构可预测存取模式油气勘探、金融分析、医疗成像、有限元、基因分析、地理信息系统、,GPGPU(GeneralPurposeComputingonGPU),GPGPU,核心思想用图形语言描述通用计算问题把数据映射到vertex或者fragment处理器缺点硬件资源使用不充分存储器访问方式严重受限难以调试和查错高度图形处理和编程技巧,CUDA(ComputeUnifiedDeviceArchitecture),CUDA有效结合CPU+GPU编程串行部分在CPU上运行并行部分在GPU上运行,CPUSerialCode,Grid0,GPUParallelKernelKernelA(args);,Grid1,CPUSerialCode,GPUParallelKernelKernelB(args);,CUDA极大提高了现有应用的效果,Gridding1,FFT,CartesianScanData,(a),(b),(b),IterativeReconstruction,(c),SpiralScanData,Spiralscandata+Gridding+FFTReconstructionrequireslittlecomputation,BasedonFig1ofLustigetal,FastSpiralFourierTransformforIterativeMRImageReconstruction,IEEEIntlSymp.onBiomedicalImaging,2004,MRIReconstruction,AdvancedMRIReconstruction,FFT,CartesianScanData,(a),SpiralScanData,IterativeReconstruction,(c),Gridding,(b),(b),Spiralscandata+IterativereconReconstructionrequiresalotofcomputation,AdvancedMRIReconstruction,ComputeQ,AcquireData,ComputeFHd,Find,Morethan99.5%oftime,Q只和扫描参数有关FHd是数据相关的使用线性求解器计算,Haldar,etal,“Anatomically-constrainedreconstructionfromnoisydata,”MRinMedicine.,Code,for(p=0;pnumP;p+)for(d=0;dnumD;d+)exp=2*PI*(kxd*xp+kyd*yp+kzd*zp);cArg=cos(exp);sArg=sin(exp);rFhDp+=rRhod*cArgiRhod*sArg;iFhDp+=iRhod*cArg+rRhod*sArg;,_global_voidcmpFhD(float*gx,gy,gz,grFhD,giFhD)intp=blockIdx.x*THREADS_PB+threadIdx.x;/registerallocateimage-spaceinputs,CPU,GPU,性能提升情况,S.S.Stone,etal,“AcceleratingAdvancedMRIReconstructionusingGPUs,”ACMComputingFrontierConference2008,Italy,May2008.,计算结果对比,CUDA成功案例,广泛应用于生命科学、机械、石油、金融、数学、天文和通信等行业,MRI(磁共振成像)GRAPPA自动校准加速网格化快速重建ComputedTomography(CT)GEDigisensSnapCT,Stone,UIUC,Batenburg,Sijbersetal,医疗成像,量子化学,KYasuda,NagoyaU,Japan,双电子积分,RI-MP2correlationenergyinQ-Chem3.1,LeslieVogt,Harvard,现有的分子动力学软件NAMD/VMD(alpharelease)GROMACS(alpharelease)HOOMDOpenMM:
分子建模https:
/simtk.org/home/openmm,分子动力学,MonteCalo模拟投资组合优化期权及衍生品定价对冲基金风险分析,CUDA中的随机数发生器,SciFinance的MonteCalo定价模型SciCompCo.,金融,序列对比蛋白质对接生物系统的随机仿真(SSA)人体视觉皮层的自组织计算模型分析基因表达的DNA微阵列工具,Schatzetal,UMaryland,生物信息学和生命科学,3DLattice-Boltzman解算器基于Lattice-Boltzman的PDE解算器用于照明的LatticeBoltzmanNavier-Stokes解算器等离子体湍流建模,ThibaultandSenocak,TolkeandKrafczy,流体动力学,GPMAD:
离子束动力学模拟FDTD法进行的光散射模拟Acceleware的解算器,FDTD加速Accelerware,电磁学和电磁力学,天气研究与预测模型(WRF)25%30%的性能提升海啸模拟,天气,大气,海洋科学与空间建模,加密编码,模式匹配,CUDA编程模型,CUDA设备与线程,计算设备(device)作为CPU(host)的协处理器有独立的存储设备(devicememory)同时启动大量线程计算密集部分使用大量线程并行的kernelGPU与CPU线程的区别GPU的线程非常轻量,线程切换1cycle,而CPU需要1000cycleGPU上的线程数1000时才能有效利用GPU的计算能力,StreamingProcessor(SP),Afullypipelined,single-issue,inordermicroprocessor2ALUsandaFPURegisterfile32-bitscalarprocessingNoinstructionfetchandschedulingNocache,StreamingMultiprocessor(SM),AnarrayofSPs8streamingprocessor2SpecialFunctionUnits(SFU)A16KBread/writesharedmemoryNotacacheButasoftware-manageddatastoreMultithreadingissuingunitInstructionandconstantcache,CUDA程序基本结构,串行部分在CPU上运行(host)并行部分在GPU上运行(device),CPUSerialCode(host),Grid0,GPUParallelKernel(device)KernelA(args);,Grid1,CPUSerialCode(host),GPUParallelKernel(device)KernelB(args);,C扩展,Declspecsglobal,device,shared,local,constantKeywordsthreadIdx,blockIdxIntrinsics_syncthreadsRuntimeAPIMemory,symbol,executionmanagementFunctionlaunch,_device_floatfilterN;_global_voidconvolve(float*image)_shared_floatregionM;.regionthreadIdx=imagei;_syncthreads().imagej=result;/AllocateGPUmemoryvoid*myimage=cudaMalloc(bytes)/100blocks,10threadsperblockconvolve(myimage);,CUDA程序的编译,使用nvcc编译工具nvcc.cu-oexcutable调试选项:
-g(debug)、-deviceemu(CPU模拟GPU),并行线程组织,并行性的维度一维y=a+b二维P=MN三维CTorMRI,并行线程组织结构,Thread:
并行的基本单位Threadblock:
互相合作的线程组CooperativeThreadArray(CTA)允许彼此同步通过快速共享内存交换数据以1维、2维或3维组织最多包含512个线程Grid:
一组threadblock以1维、2维或3维组织共享全局内存Kernel:
在GPU上执行的核心程序Onekernelonegrid,线程层次,BlockandThreadIDs,Blocks和Threads具有IDsthreadIdx,blockIdxBlockID:
1Dor2DThreadID:
1D,2Dor3D由此决定相应处理数据,CUDA线程组织,CUDAkernel函数由一系列线程组成单指令多数据流(SPMD)通过IDs确定处理的数据线程可划分为不同的Block在同一个block中,可以通过sharememory、atomicoperation和barriersynchronization进行协同,ThreadBlock0,ThreadBlock1,ThreadBlockN-1,一个简单的例子IncrementArrayElements,/CPUprogramvoidinc_cpu(float*a,floatb,intN)for(intidx=0;idxN;idx+)aidx=aidx+b;voidmain()inc_cpu(a,b,N);,/CUDAprogram_global_voidinc_gpu(float*a,floatb,intN)intidx=blockIdx.x*blockDim.x+threadIdx.x;if(idx(a,b,N);,CUDA线程的同步,void_syncthreads();Barriersynchronization同步threadblock之内的所有线程避免访问共享内存时发生RAW/WAR/WAW冒险,_shared_floatscratch256;scratchthreadID=beginthreadID;_syncthreads();intleft=scratchthreadID-1;,在此等待,直至所有线程到达才开始执行下面的代码,存储器模型与内存分配,R/Wper-threadregisters1-cyclelatencyR/Wper-threadlocalmemorySlowregis
- 配套讲稿:
如PPT文件的首页显示word图标,表示该PPT已包含配套word讲稿。双击word图标可打开word文档。
- 特殊限制:
部分文档作品中含有的国旗、国徽等图片,仅作为作品整体效果示例展示,禁止商用。设计者仅对作品中独创性部分享有著作权。
- 关 键 词:
- cuda 程序设计