基于GPU的并行程序设计.ppt

1、,2020年8月19日星期三,程序设计语言范型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）, Ad

2、dison Wesley, 2005. ,参考资料,第二部分 基于GPU的并行程序设计,第七章,引言,3. GPU高性能运算之CUDA 张舒,褚艳利 中国水利水电出版社 2009年10月,参考资料,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 1.1 GPU体系结构的演变 1.2 GPU编程模型的演变 2.CUDA编程模型 3.CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线 fixed-function GPU（功能固定的GPU）,顶点处理 （VP）,光栅化 （Rasterization）,片

3、元 处理 FP,像素 操作 PO,Memory Buffer,V-Vertex（顶点） F-Fragment（片元） P-Pixel（像素）,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线 fixed-function GPU（功能固定的GPU）,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线 fixed-function GPU（功能固定的GPU） 特点： 已经采用并行处理结构 任务并行 数据并行 采用功能固定的设计方式,第二部分 基于GPU的并行程序设计,第七

4、章,1.1 GPU体系结构的演变,（2）可编程的GPU结构及图形绘制流水线,顶点 生成 VG,顶点 处理 VP,图元 生成 PG,图元 处理 PP,片元 生成 FG,片元 处理 FP,像素 操作 PO,Memory Buffer,顶点描述,顶点数据缓存,全局缓存,纹理,顶点拓扑,全局缓存,纹理,全局缓存,纹理,输出图像,shader函数,V-Vertex（顶点）P-Primitive（图元） F-Fragment（片元） P-Pixel（像素）,可 编 程GPU结 构 图,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（3）Unified Shader GPU（统一

5、渲染架构） 主要目的是再提升可编程GPU的效率。 概念的由来： 由于图形的处理流程采用流水线的方式，必须要等到上一个阶段处理完成才能进行下一个阶段的工作。 这容易导致性能问题。,G80是首个支持DirectX10，使用“统一渲染架构”的显示核心,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,2007年，G80并行计算平台（G80 图形模式）,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,Unified Shader,第二部分 基于GPU的并行程序设计,第七章,1.1

6、 GPU体系结构的演变,Unified Shader GPU的图形绘制流水线,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU现状 高度并行化、多线程、高存储器带宽、众核、可编程,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,面

7、向通用计算 大量的晶体管用于Cache和控制电路,CPU,GPU,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU与CPU,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) 每个S

8、M中包含了8个流处理器SP（共128个SP） 每个SM管理了24个线程簇(warp)，共有768个线程 32个Threads/warp 12288个Threads/GPU 每个SM是一个SIMD处理器 每个周期在8个SP上并行执行一个线程簇,SM,第二部分 基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,多核CPU（multicore CPUs）+众核GPU（manycore GPUs） 目前主流的处理器芯片已成为并行系统!,如何开发可透明地扩展并行性的应用软件， 以便利用数量日益增加的处理器内核？,Thinking Parallel！,第二部分 基于GPU的并行程序设计,第七

9、章,1.1 GPU体系结构的演变,基于CUDA的并行应用开发实例 应用领域 计算生物学、物理模拟、图像视频处理 经济社会、气象预报、航空航天、集装箱检查 油气勘探、离散模拟、三维扫描、三维CT重建 基因、蛋白质和DNA微观生命科学研究 适用于： 对多个数据进行同一种运算（STMD）- 数据级并行 一次存储器访问，多次运算（外部DDR访问开销高，局部存储器容量较小） 浮点计算比例高（特别是单精度浮点） 今后发展方向：通用计算（Nvidia Tesla架构） 任务级并行+数据级并行+双精度浮点计算,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 1.1 GPU体系结构的演变 1.2 G

10、PU编程模型的演变 2.CUDA编程模型 3.CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,（1）GPGPU 可编程GPU的结构已经具有很高的运算能力！ 程序员对程序的行为有更多的控制能力。 增加了GPU内部结构ALU（算术逻辑单元）的数量 导致了General Purpose GPU概念的产生，即GPGPU。,GPU与CPU的运算比较,（ General-Purpose Computing on GPUs ）,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPGPU程序运行架构：,第二部分 基于GPU的并行程序设计

11、,第七章,1.2 GPU编程模型的演变,GPGPU的缺点 必须使用shader语言编写程序（学习周期较长）。 从结构图来看，CPU到GPU之间的数据传输路径较远，而且必须重复在GPU和DRAM之间进行数据的传输，效率低下。,GPU编程模型的挑战： 进一步提高GPU的高性能和通用性。 设计“GPU计算接口”，“抽象”GPU的并行计算，达到易用性。 设计基于GPU的并行计算架构及软件开发环境。,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,（2）CUDA编程模型 CUDA:Compute Unified Device Architecture 统一设备计算架构：一种通用

12、并行计算架构 2007 推出第一代完全支持CUDA的GPU产品-G80 2008推出第二代完全支持CUDA的GPU产品-GT200,Nvidia GPU (支持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语言的一些扩展 不需要计算机图形学的知识

13、 （2）没有图形API的额外开销 （3）无限制、随机访问字节内存 线程可访问内存任何地方 线程可根据需要读写多个内存位置 （4）提供了便于用户使用的并行抽象方法 编程模型 存储器模型 通信机制,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPU硬件和CUDA软件安装: ,安装顺序(注：Nvidia 8系列以上的显卡支持CUDA） CUDA Driver CUDA Toolkit CUDA SDK code samples,第二部分 基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,CUDA软件开发工具包,第二部分 基于GPU的并行程序设计,第七章,1

14、.2 GPU编程模型的演变,NVCC,C/C+ CUDA Application,PTX to Target Compiler,G80,GPU,PTX Code,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

15、.4 存储器的层次结构 2.5 CUDA 变量类型修饰符 2.6 CUDA 设备存储器分配 2.7 CUDA Host-Device 数据传输 2.8 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.1 CUDA编程接口,包括： C语言的最小扩展集(使源代码的某些部分可在设备上执行) 函数类型限定符；四个内置变量 变量类型限定符；一条新指令 运行时库 主机组件 运行在主机上，提供函数来通过主机控制和访问一个或多个计算设备。 设备组件 运行在设备上，提供特定于设备的函数。 通用组件 提供内置向量类型和C标准库的一个子集，主机和设备代码中都将支持此子集

16、。,第二部分 基于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 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.2 一个简单的CUDA程序,【例1】长度为N的向量A和B求和，结果存储在C中。,_global_ void vecAdd(float* A, float* B, float* C) i

17、nt 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),串行代码,并行内核函数2 KernelB(args),第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 2.1 CUDA编程接口 2.2 一个简单的CUDA程序 2.3 线程的层

18、次结构 2.4 存储器的层次结构 2.5 CUDA 变量类型修饰符 2.6 CUDA 设备存储器分配 2.7 CUDA Host-Device数据传输 2.8 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.3 线程的层次结构,线程 CUDA中的基本执行单元； 硬件支持，开销很小； 所有线程执行相同的代码 SIMT：单指令多线程 SPMD：单程序多数据 每个线程有一个唯一的标识ID（threadIdx） 若干线程可以组成块(Block) Thread ID: 1D、2D或3D 若干块可以组成网格(Grid) Block ID: 1D或2D,第二部分

19、 基于GPU的并行程序设计,第七章,2.3 线程的层次结构,C语言扩展:四个内置变量 dim3 gridDim; Grid的维度 (gridDim.z unused) uint3 blockIdx; Grid中Block的索引 dim3 blockDim; Block的维度 uint3 threadIdx; Block中Thread的索引,限制条件： 不允许提取任何内置变量的地址，不允许为任何内置变量赋值！,第二部分 基于GPU的并行程序设计,第七章,2.3 线程的层次结构,【例2】 长度为N（N=100000）的向量A和B求和，结果存储在C中,Block（0）,Block（1）,Block（

20、2）,Block（390）,Grid,Thread（0） Thread（1） . Thread（255）,【例3】 两个N*N(N=1024)的矩阵A和B求和，将结果存储在矩阵C中,_global_ void MatAdd(float A, float B, float C) int i = blockDim.x * blockIdx.x + threadIdx.x; /列 int j = blockDim.y * blockIdx.y + threadIdx.y; /行 if (i N ,int main() dim3 dimBlock(16 16); dim3 dimGrid (N + d

21、imBlock.x-1)/ dimBlock.x, (N + dimBlock.y-1)/ dimBlock.y); MatAdd(A, B, C); ,G80 示例: 执行线程块,Blocks,Blocks,SM1,SM0,.,SM15,8个标量处理器（SP）核心 2个特殊函数单元（SFU） 1个多线程指令单元（MTIU） MTIU：Multi-threaded Instruction Fecth and Issue Unit 16K芯片片上共享存储器(Shared Memory),【说明1】（以G80为例） Assignment：线程被指派到SM的方式。 执行kernel时，线程被组织成b

22、lock放入其中一个SM中。 每个SM最多放8个Blocks；每个Block最多512个线程。 G80中的SM最多可以有768个线程。 可以是256 (threads/block) * 3 blocks 或者是128 (threads/block) * 6 blocks, 等等。,思考： 假定两个1024*1024的矩阵相加 256（16*16）个Threads/Block 64*64个blocks要处理 每次放入SM多少个Blocks？ 每次多少个Blocks同时运行？ 每次多少个Threads并行执行？,3,3*16=48,48*256=12288,【说明2】（以G80为例） Thread

23、 Scheduling：线程调度 每个线程Block以32个线程作为一个线程簇（Warp）执行。 这是实现策略，而不是CUDA编程模型的一部分。 Warp是SM中的调度单元。,思考： 假定一个SM指派3个blocks，并且256个Threads/Block， 那么在一个SM中有多少个 每个Block被分割成256/32=8个Warps 每个SM共3*8=24个Warps。 在任何时间点，24个Warps只有其中的一个在,Block 1 Warps,Block 2 Warps,SP,SP,SP,SP,SFU,SP,SP,SP,SP,SFU,Instruction Fetch/Dispatch,I

24、nstruction L1,Streaming Multiprocessor,Shared Memory,Warps？,硬件中执行。,【说明2】（以G80为例）（续） Thread Scheduling：线程调度 SM硬件实现Zero-Overhead的Warp（32个连续的线程）调度！ 在G80中，对于所有的线程执行相同的指令，每条指令需要4个clock cycles。,【说明3】 Thread Synchronization：线程同步 一个块内的线程可彼此协作，通过共享存储器来共享数据，并栅栏同步（barrier synchronization）其执行来协调存储器的访问。 通过调用_syn

25、cthreads()内置函数在内核中指定同步点 void _syncthreads(); 位于不同block中的线程无法彼此协作! 一个线程块的线程在同一个SM上并发执行。 在线程块终止时，将在空闲的多处理器上启动新块。 SM会在硬件中创建、管理和执行并发线程，而调度开销保持为0。 CUDA关键特性： Block中的线程间的“细粒度并行” 快速的barrier同步+轻量级线程创建+零开销线程调度 Block间的“粗粒度并行”,【例4】线程同步 使用shared memory和同步求一个32x128矩阵中的每行之和，并将其存储到向量B的对应元素中。,_global_ void Sum(float

26、 *A, float *B) int mian( ) int bid=blockIdx.x; / 调用kernel int tid=threadIdx.x; Sum(A, B); _shared_ s_data128; / 将数据从global memory读入shared memory s_data128=Abid*128+tid; _syncthreads(); / 归约求和 for(int i=64; i0; i/=2) if(tidi) s_datatid=s_datatid+s_datatid+i; _syncthreads(); if(tid=0) Btid=s_data0; ,【

27、说明4】 Transparent Scalability：透明的扩展性 Hardware is free to assigns blocks to any processor at any time. A kernel scales across any number of parallel processors,Each block can execute in any order relative to other blocks.,time,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 2.1 CUDA编程接口 2.2 一个简单的CUDA程序 2.3 线

28、程的层次结构 2.4 存储器的层次结构 2.5 CUDA 变量类型修饰符 2.6 CUDA 设备存储器分配 2.7 CUDA Host-Device 数据传输 2.8 CUDA的函数声明 3 CUDA并行编程示例,片上存储器 （On-Chip Memory）,板载显存 （On-Chip Memory）,硬件模型,片上存储器 （On-Chip Memory）,第二部分 基于GPU的并行程序设计,第七章,2.4 存储器的层次结构,存储器的层次结构 CUDA线程在执行过程中可访问多个存储器空间的数据,Per-thread local memory (Read/Write),Per-block sha

29、red memory (Read/Write),Global Memory （Read/write） Constant Memory （Read） Texture Memory （Read）,第二部分 基于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 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.

30、5 CUDA 变量类型修饰符,CUDA 变量类型修饰符 当与_local_, _shared_或_constant_一起使用时， _device_可以省略。 没有限定符的自动变量默认驻留在register中。 数组除外，它们将会被放在local memory中。,第二部分 基于GPU的并行程序设计,第七章,2.5 CUDA 变量类型修饰符,CUDA 变量类型修饰符（续) 在哪里声明变量?,主机能否访问?,在任何函数外,在内核,可以,不可以,global constant,register (automatic) shared local,第二部分 基于GPU的并行程序设计,第七章,内容,1.引

31、言 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 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.6 CUDA设备存储器分配,cudaMalloc() 在全局存储器中分配空间 两个参数： 地址指针 空间大小 cudaFree() 在全局存储器中回收空间 一个参数 回收空间地址指针,第二部分 基于GPU的并行程序设计,第七章,2.6 CUDA设备

32、存储器分配,示范代码 分配64 * 64的单精度浮点数组 存储器地址为Md.elements,TILE_WIDTH = 64; Matrix Md /*“d” is often used to indicate a device data structure*/ int size = TILE_WIDTH * TILE_WIDTH * sizeof(float); cudaMalloc(void*),第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 2.1 CUDA编程接口 2.2 一个简单的CUDA程序 2.3 线程的层次结构 2.4 存储器的层次结构 2.5

33、 CUDA变量类型修饰符 2.6 CUDA设备存储器分配 2.7 CUDA Host-Device数据传输 2.8 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.7 CUDA Host-Device 数据传输,cudaMemcpy() 存储器数据传输 参数： 目的地址 源地址 传输字节数 传输类型 设备设备 主机全局存储器 全局存储器主机 异步传输,第二部分 基于GPU的并行程序设计,第七章,2.7 CUDA Host-Device 数据传输,示范代码 传输 64 * 64 单精度浮点数组 M在主机存储器中 Md在设备存储器中 传输方向常数： c

34、udaMemcpyHostToDevice cudaMemcpyDeviceToHost,cudaMemcpy(Md.elements, M.elements, size, cudaMemcpyHostToDevice); cudaMemcpy(M.elements, Md.elements, size, cudaMemcpyDeviceToHost);,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 2.1 CUDA编程接口 2.2 一个简单的CUDA程序 2.3 线程的层次结构 2.4 存储器的层次结构 2.5 CUDA变量类型修饰符 2.6 CUDA设备

35、存储器分配 2.7 CUDA Host-Device数据传输 2.8 CUDA的函数声明 3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计,第七章,2.8 CUDA的函数声明,CUDA的函数声明 _global_定义一个内核函数 必须返回void _device_ 和 _host_可以一起使用,第二部分 基于GPU的并行程序设计,第七章,2.8 CUDA的函数声明,CUDA的函数声明(续) _device_ 函数不能使用函数指针； _device_ 函数在设备上执行，所以: 不能递归 不能在函数内定义静态变量 No variable number of arguments,调用一个

36、内核函数线程生成,调用内核函数必须使用执行配置(execution configuration):,_global_ void KernelFunc(.); dim3 DimGrid(100, 50); / 5000 thread blocks dim3 DimBlock(4, 8, 8); / 256 threads per block size_t SharedMemBytes=64; / 每个线程需要64字节共享内存,KernelFunc(.);,/ 用指明调用内核函数的执行配置,任何从CUDA 1.0调用的内核函数都是异步的, 否则需要显式的同步。,第二部分 基于GPU的并行程序设计,第七章,内容,1.引言 2.CUDA编程模型 3 CUDA并行编程示例 3.1 通用编程策略 3.2 CUDA编程框架 3.3 矩阵相乘（使用全局内存） 3.4 矩阵相乘（使用共享内存）,第二部分 基于GPU的并行程序设计,第七章,3