《第七章基于GPU的并行程序设计-2教材课程》由会员分享，可在线阅读，更多相关《第七章基于GPU的并行程序设计-2教材课程（80页珍藏版）》请在金锄头文库上搜索。
1、,2020/8/10,程序设计语言范型Programming Languages Paradigms,教师： 张荣华 华北电力大学计算机系软件教研室（保定）,第七章 并行程序设计范型,基于GPU的并行程序设计,第二部分,第二部分 基于GPU的并行程序设计,第七章,引言,1.Nvidia, CUDA Programmng Guide（CUDA 2.0） 2.M. Pharr (ed.), GPU Gems 2 （Programming Techniques for High Performance Graphics and General-Purpose Computation）, Addiso
2、n Wesley, 2005. ,参考资料,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 1.1 GPU体系结构的演变 1.2 GPU编程模型的演变 2.CUDA编程模型 3.CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线 fixed-function GPU（功能固定的GPU）,顶点处理 （VP）,光栅化 （Rasterization）,片元 处理 FP,像素 操作 PO,Memory Buffer,V-Vertex（顶点） F-Fragment（片元） P-Pixel（像素）,第二部分 基
3、于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线 fixed-function GPU（功能固定的GPU）,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线 fixed-function GPU（功能固定的GPU） 特点： 已经采用并行处理结构 任务并行 数据并行 采用功能固定的设计方式,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（2）可编程的GPU结构及图形绘制流水线,顶点 生成 VG,顶点 处理 VP,图元 生成 PG,图元 处理 PP,片元
4、生成 FG,片元 处理 FP,像素 操作 PO,Memory Buffer,顶点描述,顶点数据缓存,全局缓存,纹理,顶点拓扑,全局缓存,纹理,全局缓存,纹理,输出图像,shader函数,V-Vertex（顶点）P-Primitive（图元） F-Fragment（片元） P-Pixel（像素）,可 编 程GPU结 构 图,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（3）Unified Shader GPU（统一渲染架构） 主要目的是再提升可编程GPU的效率。 概念的由来： 由于图形的处理流程采用流水线的方式，必须要等到上一个阶段处理完成才能进行下一个阶段的工作。
5、 这容易导致性能问题。,G80是首个支持DirectX10，使用“统一渲染架构”的显示核心,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,2007年，G80并行计算平台（G80 图形模式）,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,Unified Shader,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,Unified Shader GPU的图形绘制流水线,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,G
6、PU现状 高度并行化、多线程、高存储器带宽、众核、可编程,G80 GeForce 8800 GTX G92 GeForce 9800 GTX GT200 GeForce GTX 280,CPU与GPU的峰值浮点计算能力比较,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU与CPU,GPU(Graphics Process Unit),面向计算密集型和大量数据并行化的计算 大量的晶体管用于数据处理,通用CPU,面向通用计算 大量的晶体管用于Cache和控制电路,CPU,GPU,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU与CPU
7、,GT200（集成了14亿个晶体管） （240个Streaming Processor或Shader core),第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GeForce 8800 GTX（G80 CUDA计算模式）,SM（流多处理器）,SP （流处理器）,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GeForce 8800 GTX（技术参数）,共有16个流多处理器(SM) 每个SM中包含了8个流处理器SP（共128个SP） 每个SM管理了24个线程簇(warp)，共有768个线程 32个Threads/warp 12288个Thr
8、eads/GPU 每个SM是一个SIMD处理器 每个周期在8个SP上并行执行一个线程簇,SM,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,多核CPU（multicore CPUs）+众核GPU（manycore GPUs） 目前主流的处理器芯片已成为并行系统!,如何开发可透明地扩展并行性的应用软件， 以便利用数量日益增加的处理器内核？,Thinking Parallel！,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,基于CUDA的并行应用开发实例 应用领域 计算生物学、物理模拟、图像视频处理 经济社会、气象预报、航空航天、集装箱检查
9、油气勘探、离散模拟、三维扫描、三维CT重建 基因、蛋白质和DNA微观生命科学研究 适用于： 对多个数据进行同一种运算（STMD）- 数据级并行 一次存储器访问，多次运算（外部DDR访问开销高，局部存储器容量较小） 浮点计算比例高（特别是单精度浮点） 今后发展方向：通用计算（Nvidia Tesla架构） 任务级并行+数据级并行+双精度浮点计算,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 1.1 GPU体系结构的演变 1.2 GPU编程模型的演变 2.CUDA编程模型 3.CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,（1）GPG
10、PU 可编程GPU的结构已经具有很高的运算能力！ 程序员对程序的行为有更多的控制能力。 增加了GPU内部结构ALU（算术逻辑单元）的数量 导致了General Purpose GPU概念的产生，即GPGPU。,GPU与CPU的运算比较,（ General-Purpose Computing on GPUs ）,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPGPU程序运行架构：,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPGPU的缺点 必须使用shader语言编写程序（学习周期较长）。 从结构图来看，CPU到GPU之间的数据传输路
11、径较远，而且必须重复在GPU和DRAM之间进行数据的传输，效率低下。,GPU编程模型的挑战： 进一步提高GPU的高性能和通用性。 设计“GPU计算接口”，“抽象”GPU的并行计算，达到易用性。 设计基于GPU的并行计算架构及软件开发环境。,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,（2）CUDA编程模型 CUDA:Compute Unified Device Architecture 统一设备计算架构：一种通用并行计算架构 2007 推出第一代完全支持CUDA的GPU产品-G80 2008推出第二代完全支持CUDA的GPU产品-GT200,Nvidia GPU
12、(支持CUDA并行计算架构）,C with CUDA extension,C with CUDA extension,C with CUDA extension,Fortran,OpenCL,Direct Compute,C+ Python .,GPU并行应用程序,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,CUDA相对传统GPGPU的优势： （1）较低的学习曲线 只是对C语言的一些扩展 不需要计算机图形学的知识 （2）没有图形API的额外开销 （3）无限制、随机访问字节内存 线程可访问内存任何地方 线程可根据需要读写多个内存位置 （4）提供了便于用户使用的并行抽
13、象方法 编程模型 存储器模型 通信机制,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPU硬件和CUDA软件安装: ,安装顺序(注：Nvidia 8系列以上的显卡支持CUDA） CUDA Driver CUDA Toolkit CUDA SDK code samples,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,CUDA软件开发工具包,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,NVCC,C/C+ CUDA Application,PTX to Target Compiler,G80,GPU,PTX C
14、ode,Virtual,Physical,CPU Code,float4 me = gxgtid; me.x += me.y * me.z;,ld.global.v4.f32 $f1,$f3,$f5,$f7, $r9+0; mad.f32 $f1, $f5, $f3, $f1;,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 2.1 CUDA编程接口 2.2 一个简单的CUDA程序 2.3 线程的层次结构 2.4 存储器的层次结构 2.5 CUDA 变量类型修饰符 2.6 CUDA 设备存储器分配 2.7 CUDA Host-Device 数据传输 2.8 C
15、UDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.1 CUDA编程接口,包括： C语言的最小扩展集(使源代码的某些部分可在设备上执行) 函数类型限定符；四个内置变量 变量类型限定符；一条新指令 运行时库 主机组件 运行在主机上，提供函数来通过主机控制和访问一个或多个计算设备。 设备组件 运行在设备上，提供特定于设备的函数。 通用组件 提供内置向量类型和C标准库的一个子集，主机和设备代码中都将支持此子集。,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 2.1 CUDA编程接口 2.2 一个简单的CUDA程序 2.3 线
16、程的层次结构 2.4 存储器的层次结构 2.5 CUDA 变量类型修饰符 2.6 CUDA 设备存储器分配 2.7 CUDA Host-Device 数据传输 2.8 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.2 一个简单的CUDA程序,【例1】长度为N的向量A和B求和，结果存储在C中。,_global_ void vecAdd(float* A, float* B, float* C) int i = threadIdx.x; Ci = Ai + Bi; int main() / 调用Kernel函数 vecAdd(A, B, C); ,内核（kernel） 在device上由 N个不同的 CUDA线程并行执行N 次,在host上执行的串行代码 注：只能在host端调用kernel函数生成并行线程,C程序（顺序执行）,串行代码,并行内核函数1 KernelA(args
