石油计算的曙光_地震数据处理中的高性能计算

DawningCo.Ltd.CopyrightsReserved.石油计算的曙光：地震数据处理中的高性能计算技术支持中心何牧君曙光信息产业（北京）有限公司2010.12DawningCo.Ltd.CopyrightsReserved.蝉联TOP100机器数量份额第一名2010中国TOP100中科院过程所1%国防科大1%宝德2%神威3%联想3%曙光34%DELL3%浪潮5%HP19%IBM28%DawningCo.Ltd.CopyrightsReserved.提纲曙光GPU服务器产品线1DawningCo.Ltd.CopyrightsReserved.GPGPU服务器设计的思考•CPU/GPU比例的考虑–Intel5520双芯片平台–8个PCI-Ex16slot支持1～8个GPGPU设备–8GPGPU是否合适？DawningCo.Ltd.CopyrightsReserved.GPU越多越好么？TeraChemDGEMM•4块C2050相比1块C2050加速比为2.3~3•在DGEMM中8块C2050性能却不如4块•程序的数据可重用性决定了GPU数量无法太多15.3613.2539.7547.701020304050601threadx56506threadx56501xC20504xC20504xGTX480Speedup0204060801001201threadX565012threadX56501*C20502*C20504*C20508*C2050SpeedupDawningCo.Ltd.CopyrightsReserved.丰富的GPGPU产品线单GPU服务器多GPU服务器A440-G工作站入门级产品A650r-GAMD平台支持1～2GPU设备PHPC200桌边个人高性能GPU服务器5节点/10CPU/5GPU40GbInfiniband/20TB存储/UPSW580I支持1～4GPU设备TC3600刀片服务器具有极佳扩展性曙光“星云”适合多种复杂GPU应用W780I支持1～8GPU设备DawningCo.Ltd.CopyrightsReserved.提纲曙光GPU服务器产品线12GPGPU应用优化DawningCo.Ltd.CopyrightsReserved.Fermi癿重要特点•双发射流多处理器•新增L1cache•Sharedmemory容量增加•双精度计算性能大幅提升，为单精度癿一半•显存ECC支持L2CacheMemoryControllerGPCSMRasterEnginePolymorphEngineSMPolymorphEngineSMPolymorphEngineSMPolymorphEngineGPCSMRasterEnginePolymorphEngineSMPolymorphEngineSMPolymorphEngineSMMemoryControllerMemoryControllerMemoryControllerMemoryControllerMemoryControllerGPCSMRasterEnginePolymorphEngineSMPolymorphEngineSMPolymorphEngineSMPolymorphEngineGPCSMRasterEnginePolymorphEngineSMPolymorphEngineSMPolymorphEngineSMPolymorphEnginePolymorphEngineHostInterfaceGigaThreadEngineGT200GF100L1TextureCache(perquad)12KB12KBDedicatedL1LD/STCacheX16or48KBTotalSharedMemory16KB16or48KBL2Cache256KB(TEXreadonly)768KB(allclientsread/write)DoublePrecisionThroughput30FMAs/clock256FMAs/clockDawningCo.Ltd.CopyrightsReserved.性能因素•Massiveparallelism•Occupancy(RegistersandSharedMemory)•HeterogeneousCPU+GPUcomputing•Control-FlowDivergence•MemoryCoalescing•SharedMemoryBankConflicts•GPUdirectDawningCo.Ltd.CopyrightsReserved.Massiveparallelism•使用成百上千个线程块–一个线程块在一个SM上执行–每个SM能够有效执行2-8个线程块–粗粒度任务级并行•每个线程块需包含数百个线程–每个SM需要192-1536个线程利用所有核心和隐藏长延迟指令–细粒度数据并行、指令并行DawningCo.Ltd.CopyrightsReserved.线程块配置选取例子•分块矩阵乘，每个SM最多能支持1536个线程（丌考虑其他资源限制），使用8x8、16x16、32x32、64x64线程块？–8x8:每个线程块64个线程，每个SM有1536/64=24个线程块，但每个最多支持8个线程块，因此只有512个线程执行核心利用率33%，而且丌利于隐藏长延迟操作。–16x16:每个线程块256个线程，每个SM有1536/256=6个线程块，核心利用率100%。–32x32:每个线程块1024个线程，每个SM有1个线程块，核心利用率67%，而且丌利于隐藏长延迟操作。–64x64:每个线程块4096个线程，没有一个线程块能再SM上执行。DawningCo.Ltd.CopyrightsReserved.OCCUPANCY•为什么需要线程调度？–多数情况下，globalmemory访问导致流水线停顿•SM实现了零开销线程调度器–任何时刻SM只能执行一个warp–根据优先级选择待执行癿•当某个warp停顿时，调度器选择其他能够执行癿warp执行，隐藏延迟开销•因此，需要最大化利用率！–每个线程寄存器个数–每个线程块使用sharedmemory数量TB1W1TB=ThreadBlock,W=WarpTB2W1TB3W1TB2W1TB1W1TB3W2TB1W2TB1W3TB3W2TimeTB1,W1stallTB3,W2stallTB2,W1stallInstruction:1234561212341278121234DawningCo.Ltd.CopyrightsReserved.资源限制TB0Registers每个SM上癿线程共享有限数量癿寄存器和sharedmemeory集合(32Kregistersand64KBL1+sharedmemory)每个线程块占有自己需要癿寄存器和sharedmemory如果其中一项资源被完全占有，丌能调度更多线程到该SM上SharedMemoryTB1TB2TB0TB1TB2TB0RegistersTB1TB0TB1SharedMemoryDawningCo.Ltd.CopyrightsReserved.HETEROGENEOUSCPU+GPUCOMPUTING•pcie2带宽10GB/s•GPUload/store带宽150GB/s•尽量减少数据传输次数•实现计算和通信重叠DeviceMemoryPCIeBridgeCPUHostMemorycudaMemcpy()DawningCo.Ltd.CopyrightsReserved.计算和通信重叠cudaMemcpyAsync()启劢异步数据传输cudaMallocHost()分配主机内存CPUGPUandGPUCPUdatatransfersOverlapwithCPUandGPUprocessing隐藏传输开销流水线:Kernel0Kernel1Kernel2Kernel3CPUCPUCPUCPUcpy=cpy=cpy=cpy=GPUGPUGPUGPUcpy=cpy=cpy=cpy=双DMADawningCo.Ltd.CopyrightsReserved.优化实例——矩阵乘•每个线程块256个线程•每个线程块共享32KBsharedmemory•每个线程使用63个寄存器•每个SM同时运行两个线程块–32×2=64KBsharedmemory–63×256×2=3225632768registersDawning[e-mail].特定领域GPU专用机配置物理化学材料计算专用机生物信息学计算专用机PHPC20040xCPU处理单元5xGPU加速单元120GB内存20TB本地存储Linux操作系统Gridview2.5GPU专用版生物电镜（EMAN）分子动力学NAMDVMDGromacsLammpsHOOMD第一性原理BigDFTTeraChem蛋白质对接和分子叠加（HEXDocking）序列纠错（CUDA-EC）蛋白质数据库检索CUDASW++GPU-HmmerMUMmerGPUDawningCo.Ltd.CopyrightsReserved.提纲曙光GPU服务器产品线123GPGPU应用优化深入的项目合作DawningCo.Ltd.CopyrightsReserved.深入的项目合作——地震数据处理•基于合成波地震成像中波场延拓–通过对计算任务进行分割，劢态负载均衡–结合cufft/fftw，并进行封装–利用CPU/GPU计算能力DawningCo.Ltd.CopyrightsReserved.深入的项目合作——其他领域•粒子碰撞模拟•ATIHD4870x2•加速比：～17x•稀疏矩阵求解•NvidiaC2050•加速比：～10x•生物电镜——EMAN•NvidiaC2050•加速比：～50x•蛋白质分析•MS-Alignment算法•NvidiaC2050•加速比：～100x05101520CSRDIAELLCOOHYBGFLOP/s星云SpMV测试整体加速效果classesbymraclassesalignallmake3dtotal40*12coresCPU13min1min48min62min40*FermiGPU7min1min5min13minDawningCo.Ltd.CopyrightsReserved.曙光GPGPU领域成功案例2009年,曙光开发成功GHPC1000年份单位节点机型GPU2008中科院过程所42A620r-TAMDRadeon4870x22009中科院地质与地球所54A620r-TNvidiaGTX295/AMDRadeon4870x22009中科院高能物理所31W580iAMDRadeon4870x2/NvidiaC10602009紫金山天文台120定制机型NvidiaGTX295/C10602009中科院网络中心45A620r-TAMDRadeon4870x22009沈阳金属所120A620r-TNvidiaGTX295/C10602009湖南师范大学10A620r-TNvidiaC10602009中国海洋大学20A620r-TNvidiaC10602009华东理工大学28A620r-TAMDRadeon58702009甘肃省超算中心16A620r-TAMDRadeon4870x22009/2010北京计算中心98TC3600AMDRadeon4870x2/NvidiaC20502010中石化物探院24W580INvidiaC20502010浙江大学计算机系19W580INvidiaC1060/C2050/GTX4802010中山大学54W780INvidiaC20502010成都超级计算中心