GPU入门技术介绍

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编程
• CUDA • OpenCL
OpenCL(全称Open Computing Language，开放 运算语言)是第一个面向异构系统通用目的并行编 程的开放式、免费标准，也是一个统一的编程环 境，便于软件开发人员为高性能计算服务器、桌 面计算系统、手持设备编写高效轻便的代码，而 且广泛适用于多核心处理器(CPU)、图形处理器 (GPU)。
CUDA编程模型
• CUDA的基本思想是支持大量的线程级并行 （Thread Level Parallel），并在硬件中动态地调 度和执行这些线程。 • 异构思想 Host + Device(Coprocessor) 让GPU来运行一些能够被高度线程化的代码。 GPU与CPU协同工作，GPU只有在计算高度数据 并行任务时才发挥作用。
GPU简介
型号 适用用户
GeForce Quadro
Tesla Tegra Ion
家庭和企业的娱乐应用，面向游戏用户 应用于图形工作站，面向专业级用户
用于高性能通用计算，面向研究人员 适用于移动设备 上网本
GPU研究现状
• NVIDIA在1999年推出了第一款GPU产品－ GeForce 256。主要任务是进行图形渲染任务， 缓解CPU压力。 • 从GPU诞生那天开始，其发展脚步就没有停止下 来，由于其独特的体系架构和超强的浮点运算能 力，人们希望将某些通用计算问题移植到GPU上 来完成以提升效率，出现了所谓的GPGPU （General Purpose Graphic Process Unit），但 是由于其开发难度较大，没有被广泛接受。 • 2006年NVIDIA推出了第一款基于Tesla架构的 GPU（G80），GPU已经不仅仅局限于图形渲染， 开始正式向通用计算领域迈进。
GPห้องสมุดไป่ตู้内部架构
• 由于CPU和GPU设计目标的不同导致了两者在架构、并 行层次和性能方面差异较大: CPU的重线程与GPU的轻线程 CPU的MIMD多核与GPU的SIMT众核（x7560） CPU内存、缓存与GPU存储器 • GPU是以大量线程实现面向吞吐量的数据并行计算，适合 于处理计算密度高、逻辑分支简单的大规模数据并行负载； 而CPU则有复杂的控制逻辑和大容量的缓存减小延迟，擅 长复杂逻辑运算。
CUDA编程模型
grid之间通过global memory交换数据 block之间不能相互通信,只能通过global memory共享数据,不要让多个block写同一 区段内容（不保证数据一致性和顺序一致 性） 同一block内的thread可以通过shared memory和同步实现通信block间粗粒度并行， block内thread细粒度并行
CUDA存储器模型
• 执行参数与内建变量的作用 各个thread和block之间的唯一不同就是 threadID和BlockID,通过内建变量控制各个 线程处理的指令和数据 CPU运行核函数时的执行参数确定GPU在 SPA上分配多少个block，在SM上分配多少 个thread
CUDA存储器模型
CUDA程序优化
active block
每个SM最多可以有768（G8x,G9x）或者 1024(GT200)个active thread 这些active thread最多可以属于8个block 还有受到SM中shared memory和register的制约 最后的active block数量是由以上四个条件中的 “短板”决定
CUDA存储器模型
• • • • • • • • Register Local Shared Global Constant Texture Host memory Pinned host memory
CUDA存储器模型
• CUDA对C的扩展：kernel执行参数 <<< >>>运算符，用来传递一些kernel执行 参数 Grid的大小和维度 Block的大小和维度 外部声明的shared memory大小 stream编号
GPU内部架构
前进
GPU内部架构
• 在GPU中，流多处理器才能被称为真正的 完整核心，整个可伸缩流处理器阵列可以 被看成是由多个流多处理器组成的多单指 令流多线程(Multiple SIMT，MSIMT)系统。 Tesla GT200架构在可编程性和灵活性与硬 件的复杂度和功耗之间取得了很好的折衷， 线程被组织成多个线程块(Thread Block)， 分配到各个流多处理器上，而每个线程块 内的线程再被以单指令流多线程的方式交 给流处理器运行。
CUDA编程模型
• 调用kernel函数时CPU 调用API将显卡端程序 的二进制代码传到 GPU • grid运行在SPA上 • block运行在SM上 • thread运行在SP上
返回
CUDA编程模型
• grid block thread Kernel不是一个完整的程序，而只是其中的 一个关键并行计算步骤 Kernel以一个网格(Grid)的形式执行，每个 网格由若干个线程块（block）组成，每一 个线程块又由最多512个线程(thread)组成。 一个grid最多可以有65535 * 65535个block 一个block总共最多可以有512个thread，在 三个维度上的最大值分别为512, 512和64
CUDA程序实例
• 矩阵乘法
矩阵乘法时间复杂度是O (abc)，其中a b c分别 表示两个矩阵大小是a× b以及b × c.
• 实验环境
CPU: Intel(R) Xeon(R) E5430 2.66GHz (8核) 内存: DDR3 1333 4G GPU: Tesla C2050 (448核) 显存: GDDR5 3G
GPU研究现状
• 2007年6月，NVIDIA推出了CUDA（Computer Unified Device Architecture计算统一设备结构）。 CUDA是一种将GPU作为数据并行计算设备的软 硬件体系。在CUDA 的架构中，不再像过去 GPGPU架构那样将通用计算映射到图形API中， 对于开发者来说，CUDA 的开发门槛大大降低了。 CUDA 的编程语言基于标准C ，因此任何有C 语 言基础的用户都很容易地开发CUDA 的应用程序。 由于这些特性，CUDA在推出后迅速发展，被广 泛应用于石油勘测、天文计算、流体力学模拟、 分子动力学仿真、生物计算、图像处理、音视频 编解码等领域。
CUDA程序实例
• 测试数据
a=b=c=1024 a=b=c=2048
• 算法
CPU单线程 CPU多线程 GPU
CUDA程序实例
• 实验结果
1024 * 1024 CPU单线程 CPU多线程 GPU 加速比 4568.300781ms 606.676025ms 17.434999ms 35 2048*2048 557428.375000ms 193791.609375ms 113.415001 1174
• CUDA driver API • CUDA runtime API
CUDA程序模板
main(){ //Allocate memory on GPU float *Md; cudaMalloc((void**)&Md, size); //Copy data from CPU to GPU cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); //Call GPU kernel function kernel<<<dimGrid, dimBlock>>> (arguments); //Copy data from GPU back to CPU CopyFromDeviceMatrix(M, Md); //Free device matrices FreeDeviceMatrix(Md); }
GPU简介
• 市场迫切需要实时、高清晰度的 3D 图形， 可编程的 GPU 已发展成为一种高度并行化、 多线程、多核心的处理器，具有杰出的计 算能力和极高的存储器带宽。 • GPU极大提升了计算机图形处理的速度、 增强了图形的质量，并促进了与计算机图 形相关其他应用领域的快速发展。与中央 处理器(Central Processing Unit，CPU)的 串行设计模式不同，GPU为图形处理设计， 具有天然的并行特性。
GPU简介
GPU，Graphics Processing Unit的简写，是 现代显卡中非常重要的一个部分，其地位与 CPU在主板上的地位一致，主要负责的任务 是加速图形处理速度。 GPU是一个高度并行化的多线程、多核心处 理器。
GPU简介
• GPU/CPU计算能力比较
GPU简介
• GPU/CPU存储器带宽比较
GPU内部架构
CPU: 强控制弱计算，更多资源用于缓存 CPU: 强计算弱控制，更多资源用于数据计算
GPU内部架构
• GPU体系架构在不断的发展，以GT200体系架构为代表对 GPU的并行层次进行分析。Tesla GT200由两部分组成， 分别是可伸缩流处理器阵列(Scalable Streaming Processor Array，SPA)和存储器系统，它们由一个片上 互联网络连接。如下图所示，可伸缩流处理器阵列由若干 个线程处理器群(Thread Processing Cluster，TPC)构成， 每个TPC包含2～3个流多处理器(Streaming Multiprocessor，SM)，每个流多处理器中包含8个流处理 器(StreamingProcessor，SP)。流处理器有独立的寄存器 和指令指针，但缺少取指和调度单元，而流多处理器才拥 有完整前端，包括取值、译码、发射等。从结构上看，每 个流多处理器相当于一个8路单指令流多数据流(Single Instruction Multiple Data，SIMD)处理器，不同的是， GPU实现了自动向量机化，NVIDIA将之命名为单指令流 多线程(Single Instruction Multiple Thread，SIMT)
CUDA编程模型
• warp warp是硬件特性带来的概念，在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完成分支的时间是多个分支 时间之和。