GPU入门技术介绍ppt课件.ppt
GPU,GPU简介GPU研究现状GPU内部架构CUDA编程模型CUDA存储器模型CUDA程序实例CUDA程序优化,GPU简介,单核处理器芯片已经到了尽头Power Wall功耗大Memory Wall存储器延迟很难降低，缓存占据70%芯片面积,GPU简介,多核和众核时代多个适当复杂度、低功耗核心并行工作时钟频率基本不变未来计算机硬件不会更快，但会更“宽”必须重新设计算法“Multicore: This is the one which will have the biggest impact on us. We have never had a problem to solve like this. A breakthrough is needed in how applications are done on multicore devices.” -Bill Gates, Microsoft,GPU简介,GPU，Graphics Processing Unit的简写，是现代显卡中非常重要的一个部分，其地位与CPU在主板上的地位一致，主要负责的任务是加速图形处理速度。GPU是一个高度并行化的多线程、多核心处理器。,GPU简介,GPU/CPU计算能力比较,GPU简介,GPU/CPU存储器带宽比较,GPU简介,市场迫切需要实时、高清晰度的 3D 图形，可编程的 GPU 已发展成为一种高度并行化、多线程、多核心的处理器，具有杰出的计算能力和极高的存储器带宽。GPU极大提升了计算机图形处理的速度、增强了图形的质量，并促进了与计算机图形相关其他应用领域的快速发展。与中央处理器(Central Processing Unit，CPU)的串行设计模式不同，GPU为图形处理设计，具有天然的并行特性。,GPU简介,GPU研究现状,NVIDIA在1999年推出了第一款GPU产品GeForce 256。主要任务是进行图形渲染任务，缓解CPU压力。从GPU诞生那天开始，其发展脚步就没有停止下来，由于其独特的体系架构和超强的浮点运算能力，人们希望将某些通用计算问题移植到GPU上来完成以提升效率，出现了所谓的GPGPU（General Purpose Graphic Process Unit），但是由于其开发难度较大，没有被广泛接受。2006年NVIDIA推出了第一款基于Tesla架构的GPU（G80），GPU已经不仅仅局限于图形渲染，开始正式向通用计算领域迈进。,GPU研究现状,2007年6月，NVIDIA推出了CUDA（Computer Unified Device Architecture计算统一设备结构）。CUDA是一种将GPU作为数据并行计算设备的软硬件体系。在CUDA 的架构中，不再像过去GPGPU架构那样将通用计算映射到图形API中，对于开发者来说，CUDA 的开发门槛大大降低了。CUDA 的编程语言基于标准C ，因此任何有C 语言基础的用户都很容易地开发CUDA 的应用程序。由于这些特性，CUDA在推出后迅速发展，被广泛应用于石油勘测、天文计算、流体力学模拟、分子动力学仿真、生物计算、图像处理、音视频编解码等领域。,GPU内部架构,CPU: 强控制弱计算，更多资源用于缓存CPU: 强计算弱控制，更多资源用于数据计算,GPU内部架构,GPU体系架构在不断的发展，以GT200体系架构为代表对GPU的并行层次进行分析。Tesla GT200由两部分组成，分别是可伸缩流处理器阵列(Scalable Streaming Processor Array，SPA)和存储器系统，它们由一个片上互联网络连接。如下图所示，可伸缩流处理器阵列由若干个线程处理器群(Thread Processing Cluster，TPC)构成，每个TPC包含23个流多处理器(Streaming Multiprocessor，SM)，每个流多处理器中包含8个流处理器(StreamingProcessor，SP)。流处理器有独立的寄存器和指令指针，但缺少取指和调度单元，而流多处理器才拥有完整前端，包括取值、译码、发射等。从结构上看，每个流多处理器相当于一个8路单指令流多数据流(Single Instruction Multiple Data，SIMD)处理器，不同的是，GPU实现了自动向量机化，NVIDIA将之命名为单指令流多线程(Single Instruction Multiple Thread，SIMT),GPU内部架构,前进,GPU内部架构,在GPU中，流多处理器才能被称为真正的完整核心，整个可伸缩流处理器阵列可以被看成是由多个流多处理器组成的多单指令流多线程(Multiple SIMT，MSIMT)系统。Tesla GT200架构在可编程性和灵活性与硬件的复杂度和功耗之间取得了很好的折衷，线程被组织成多个线程块(Thread Block)，分配到各个流多处理器上，而每个线程块内的线程再被以单指令流多线程的方式交给流处理器运行。,GPU内部架构,由于CPU和GPU设计目标的不同导致了两者在架构、并行层次和性能方面差异较大:CPU的重线程与GPU的轻线程CPU的MIMD多核与GPU的SIMT众核（x7560）CPU内存、缓存与GPU存储器GPU是以大量线程实现面向吞吐量的数据并行计算，适合于处理计算密度高、逻辑分支简单的大规模数据并行负载；而CPU则有复杂的控制逻辑和大容量的缓存减小延迟，擅长复杂逻辑运算。,GPU编程,CUDAOpenCLOpenCL(全称Open Computing Language，开放运算语言)是第一个面向异构系统通用目的并行编程的开放式、免费标准，也是一个统一的编程环境，便于软件开发人员为高性能计算服务器、桌面计算系统、手持设备编写高效轻便的代码，而且广泛适用于多核心处理器(CPU)、图形处理器(GPU)。,CUDA编程模型,CUDA的基本思想是支持大量的线程级并行（Thread Level Parallel），并在硬件中动态地调度和执行这些线程。异构思想Host + Device(Coprocessor)让GPU来运行一些能够被高度线程化的代码。GPU与CPU协同工作，GPU只有在计算高度数据并行任务时才发挥作用。,CUDA编程模型,调用kernel函数时CPU调用API将显卡端程序的二进制代码传到GPUgrid运行在SPA上block运行在SM上thread运行在SP上,返回,CUDA编程模型,grid block threadKernel不是一个完整的程序，而只是其中的一个关键并行计算步骤Kernel以一个网格(Grid)的形式执行，每个网格由若干个线程块（block）组成，每一个线程块又由最多512个线程(thread)组成。一个grid最多可以有65535 * 65535个block一个block总共最多可以有512个thread，在三个维度上的最大值分别为512, 512和64,CUDA编程模型,grid之间通过global memory交换数据block之间不能相互通信,只能通过global memory共享数据,不要让多个block写同一区段内容（不保证数据一致性和顺序一致性）同一block内的thread可以通过shared memory和同步实现通信block间粗粒度并行，block内thread细粒度并行,CUDA编程模型,warpwarp是硬件特性带来的概念，在CUDA C语言中是透明的（除vote函数），但应用中不能忽略一个warp中有32个线程，这是因为SM中有8个SP，执行一条指令的延迟是4个周期，使用了流水线技术一个half warp中有16个线程，这是因为执行单元的频率是其他单元的两倍，每两个周期才进行一次数据传输,CUDA编程模型,分支性能与现代的微处理器不同，NVIDIA的SM没有预测执行机制-没有分支预测单元(Branch Predicator)。 在需要分支时，只有当warp中所有的线程都计算出各自的分支的地址，并且完成取指以后，warp才能继续往下执行。 如果一个warp内需要执行N个分支，那么SM就需要把每一个分支的指令发射到每一个SP上，再由SP根据线程的逻辑决定需不需要执行。这是一个串行过程，此时SIMT完成分支的时间是多个分支时间之和。,CUDA存储器模型,RegisterLocalSharedGlobalConstantTextureHost memoryPinned host memory,CUDA存储器模型,CUDA对C的扩展：kernel执行参数运算符，用来传递一些kernel执行参数 Grid的大小和维度Block的大小和维度外部声明的shared memory大小stream编号,CUDA存储器模型,执行参数与内建变量的作用各个thread和block之间的唯一不同就是threadID和BlockID,通过内建变量控制各个线程处理的指令和数据CPU运行核函数时的执行参数确定GPU在SPA上分配多少个block，在SM上分配多少个thread,CUDA存储器模型,CUDA driver APICUDA runtime API,CUDA程序模板,main()/Allocate memory on GPUfloat *Md;cudaMalloc(void*),CUDA程序实例,矩阵乘法矩阵乘法时间复杂度是O (abc)，其中a b c分别表示两个矩阵大小是a b以及b c.实验环境CPU: Intel(R) Xeon(R) E5430 2.66GHz (8核)内存: DDR3 1333 4GGPU: Tesla C2050 (448核)显存: GDDR5 3G,CUDA程序实例,测试数据a=b=c=1024a=b=c=2048算法CPU单线程CPU多线程GPU,CUDA程序实例,实验结果,CUDA程序优化,active block每个SM最多可以有768（G8x,G9x）或者1024(GT200)个active thread这些active thread最多可以属于8个block还有受到SM中shared memory和register的制约最后的active block数量是由以上四个条件中的“短板”决定,CUDA程序优化,指令优化选用计算复杂度较小的算法，处理字长为32bit并行度高，粗粒度并行多，细粒度并行有局部性分支映射成固定运算，展开代码避免循环在精度允许的前提下使用带有_前缀的快速算法，只在必要的部分使用双精度和64-bit int使用移位运算代替整数除法和求余使用vote,atomic,red等intrinsic函数实现算法只在线程间通信前进行同步,CUDA程序优化,存储器访问优化对显存的I/O成本很高，提高两次对显存访问之间的计算量，通过同时计算与访存隐藏延时把适合的数据放入纹理和常数缓存等缓解带宽压力，提高读取速度避免bank conflict，非合并访问或cpu-gpu数据传输 使用数组的结构体,而不是结构体数组使用对齐，类型转换等手段实现合并访问,
