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

《基于GPU的并行程序设计.ppt》由会员分享，可在线阅读，更多相关《基于GPU的并行程序设计.ppt（80页珍藏版）》请在三一办公上搜索。

1、,2023/10/11,程序设计语言范型Programming Languages Paradigms,教师：张荣华 华北电力大学计算机系软件教研室（保定）,第七章 并行程序设计范型,基于GPU的并行程序设计,第二部分,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,引言,1.Nvidia,CUDA Programmng Guide（CUDA 2.0）http:/2.M.Pharr(ed.),GPU Gems 2（Programming Techniques for High Performance Graphics and

2、General-Purpose Computation）,Addison Wesley,2005.（英文电子在线阅读）,参考资料,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,引言,3.GPU高性能运算之CUDA张舒,褚艳利中国水利水电出版社2009年10月,参考资料,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言1.1 GPU体系结构的演变1.2 GPU编程模型的演变2.CUDA编程模型3.CUDA并行编程示例,第二部分 基于GPU的并行程

3、序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线fixed-function GPU（功能固定的GPU）,顶点处理（VP）,光栅化（Rasterization）,片元处理 FP,像素操作 PO,Memory Buffer,V-Vertex（顶点）F-Fragment（片元）P-Pixel（像素）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线fixed

4、-function GPU（功能固定的GPU）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（1）传统的GPU结构及图形绘制流水线fixed-function GPU（功能固定的GPU）特点：已经采用并行处理结构任务并行数据并行采用功能固定的设计方式,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（2）可编程的GPU结构及图形绘制流水线,顶点生成 VG,顶点处理 VP,图元生成 PG,图元处理

5、PP,片元生成 FG,片元处理 FP,像素操作 PO,Memory Buffer,顶点描述,顶点数据缓存,全局缓存,纹理,顶点拓扑,全局缓存,纹理,全局缓存,纹理,输出图像,shader函数,V-Vertex（顶点）P-Primitive（图元）F-Fragment（片元）P-Pixel（像素）,可编程GPU结构图,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,（3）Unified Shader GPU（统一渲染架构）主要目的是再提升可编程GPU的效率。概念的由来：由于图形的处理流程采用流水线的

6、方式，必须要等到上一个阶段处理完成才能进行下一个阶段的工作。这容易导致性能问题。,G80是首个支持DirectX10，使用“统一渲染架构”的显示核心,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,2007年，G80并行计算平台（G80 图形模式）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计

7、,第七章,1.1 GPU体系结构的演变,Unified Shader,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,Unified Shader GPU的图形绘制流水线,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU现状 高度并行化、多线程、高存储器带宽、众核、可编程,G80 GeForce 8800 GTX G92 GeForce 9800 GTXGT200 GeForce GTX 280

8、,CPU与GPU的峰值浮点计算能力比较,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU与CPU,GPU(Graphics Process Unit),面向计算密集型和大量数据并行化的计算大量的晶体管用于数据处理,通用CPU,面向通用计算大量的晶体管用于Cache和控制电路,CPU,GPU,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GPU与CPU,GT200（集成了14亿个晶体管）（240

9、个Streaming Processor或Shader core),第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GeForce 8800 GTX（G80 CUDA计算模式）,SM（流多处理器）,SP（流处理器）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.1 GPU体系结构的演变,GeForce 8800 GTX（技术参数）,共有16个流多处理器(SM)每个SM中包含了8个流处理器SP（共128个SP）每个SM管理了2

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

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

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

13、e Computing on GPUs）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPGPU程序运行架构：,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPGPU的缺点必须使用shader语言编写程序（学习周期较长）。从结构图来看，CPU到GPU之间的数据传输路径较远，而且必须重复在GPU和DRAM之间进行数据的传输，效率低下。,GPU编程模型的挑战：进一步提高GPU的高性能和通用性。设计

14、“GPU计算接口”，“抽象”GPU的并行计算，达到易用性。设计基于GPU的并行计算架构及软件开发环境。,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,（2）CUDA编程模型CUDA:Compute Unified Device Architecture统一设备计算架构：一种通用并行计算架构2007 推出第一代完全支持CUDA的GPU产品-G802008推出第二代完全支持CUDA的GPU产品-GT200,Nvidia GPU(支持CUDA并行计算架构）,Cwith CUDA extension,C

15、with CUDA extension,Cwith CUDA extension,Fortran,OpenCL,DirectCompute,C+Python.,GPU并行应用程序,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,CUDA相对传统GPGPU的优势：（1）较低的学习曲线只是对C语言的一些扩展不需要计算机图形学的知识（2）没有图形API的额外开销（3）无限制、随机访问字节内存线程可访问内存任何地方线程可根据需要读写多个内存位置（4）提供了便于用户使用的并行抽象方法编程模型存储器模型通信机

16、制,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,GPU硬件和CUDA软件安装:http:/,安装顺序(注：Nvidia 8系列以上的显卡支持CUDA）CUDA Driver CUDA Toolkit CUDA SDK code samples,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,CUDA软件开发工具包,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基

17、于GPU的并行程序设计,第七章,1.2 GPU编程模型的演变,NVCC,C/C+CUDAApplication,PTX to TargetCompiler,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的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言2.CUDA编程模型2.1 CUDA

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

19、在设备上，提供特定于设备的函数。通用组件提供内置向量类型和C标准库的一个子集，主机和设备代码中都将支持此子集。,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于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的并行程序设计基于GPU的并行程序设计基于GPU的并

20、行程序设计基于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程序（顺序执行）,串行代码,并行内核函数1KernelA(args),串行代码,并行内核函数

21、2KernelB(args),第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于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的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.3 线程的层次结构,线程CUDA中的

22、基本执行单元；硬件支持，开销很小；所有线程执行相同的代码SIMT：单指令多线程SPMD：单程序多数据每个线程有一个唯一的标识ID（threadIdx）若干线程可以组成块(Block)Thread ID:1D、2D或3D 若干块可以组成网格(Grid)Block ID:1D或2D,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.3 线程的层次结构,C语言扩展:四个内置变量dim3 gridDim;Grid的维度(gridDim.z unused)uint3 blockIdx;Grid中Block的索引dim3 blockDi

23、m;Block的维度uint3 threadIdx;Block中Thread的索引,限制条件：不允许提取任何内置变量的地址，不允许为任何内置变量赋值！,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.3 线程的层次结构,【例2】长度为N（N=100000）的向量A和B求和，结果存储在C中,Block（0）,Block（1）,Block（2）,Block（390）,Grid,Thread（0）Thread（1）.Thread（255）,【例3】两个N*N(N=1024)的矩阵A和B求和，将结果存储在矩阵C中,_global_

24、 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+dimBlock.x-1)/dimBlock.x,(N+dimBlock.y-1)/dimBlock.y);MatAdd(A,B,C);,G80 示例:执行线程块,Blocks,Blocks,SM1,SM0,.,SM15,8个标量处理器（SP）核心 2

25、个特殊函数单元（SFU）1个多线程指令单元（MTIU）MTIU：Multi-threaded Instruction Fecth and Issue Unit 16K芯片片上共享存储器(Shared Memory),【说明1】（以G80为例）Assignment：线程被指派到SM的方式。执行kernel时，线程被组织成block放入其中一个SM中。每个SM最多放8个Blocks；每个Block最多512个线程。G80中的SM最多可以有768个线程。可以是256(threads/block)*3 blocks或者是128(threads/block)*6 blocks,等等。,思考：假定两个10

26、24*1024的矩阵相加 256（16*16）个Threads/Block 64*64个blocks要处理 每次放入SM多少个Blocks？每次多少个Blocks同时运行？每次多少个Threads并行执行？,3,3*16=48,48*256=12288,【说明2】（以G80为例）Thread Scheduling：线程调度 每个线程Block以32个线程作为一个线程簇（Warp）执行。这是实现策略，而不是CUDA编程模型的一部分。Warp是SM中的调度单元。,思考：假定一个SM指派3个blocks，并且256个Threads/Block，那么在一个SM中有多少个 每个Block被分割成256/

27、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,Instruction L1,Streaming Multiprocessor,Shared Memory,Warps？,硬件中执行。,【说明2】（以G80为例）（续）Thread Scheduling：线程调度SM硬件实现Zero-Overhead的Warp（32个连续的线程）调度！在G80中，对于所有的线程执行相同

28、的指令，每条指令需要4个clock cycles。,【说明3】Thread Synchronization：线程同步一个块内的线程可彼此协作，通过共享存储器来共享数据，并栅栏同步（barrier synchronization）其执行来协调存储器的访问。通过调用_syncthreads()内置函数在内核中指定同步点void _syncthreads();位于不同block中的线程无法彼此协作!一个线程块的线程在同一个SM上并发执行。在线程块终止时，将在空闲的多处理器上启动新块。SM会在硬件中创建、管理和执行并发线程，而调度开销保持为0。CUDA关键特性：Block中的线程间的“细粒度并行”快速

29、的barrier同步+轻量级线程创建+零开销线程调度Block间的“粗粒度并行”,【例4】线程同步使用shared memory和同步求一个32x128矩阵中的每行之和，并将其存储到向量B的对应元素中。,_global_ void Sum(float*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();/归约求和

30、 for(int i=64;i0;i/=2)if(tidi)s_datatid=s_datatid+s_datatid+i;_syncthreads();if(tid=0)Btid=s_data0;,【说明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 relativ

31、e to other blocks.,time,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于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并行编程示例,片上存储器（On-Chip Memory）,板载显存（On-Chip Memory）,硬件模型,片上存储器（On-Chip Memory

32、）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.4 存储器的层次结构,存储器的层次结构CUDA线程在执行过程中可访问多个存储器空间的数据,Per-thread local memory(Read/Write),Per-block shared memory(Read/Write),Global Memory（Read/write）Constant Memory（Read）Texture Memory（Read）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,

33、第七章,内容,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的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.5 CUDA 变量类型修饰符,CUDA 变量类型修饰符当与_local_,_shared_或_constant_一起使用时，_device_可以省略。没有限定符的自动变量默认驻

34、留在register中。数组除外，它们将会被放在local memory中。,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.5 CUDA 变量类型修饰符,CUDA 变量类型修饰符（续)在哪里声明变量?,主机能否访问?,在任何函数外,在内核,可以,不可以,globalconstant,register(automatic)sharedlocal,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言2.CUDA编程模型2.1 CUDA编程接口2.2

35、 一个简单的CUDA程序2.3 线程的层次结构2.4 存储器的层次结构2.5 CUDA变量类型修饰符2.6 CUDA设备存储器分配2.7 CUDA Host-Device 数据传输2.8 CUDA的函数声明3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.6 CUDA设备存储器分配,cudaMalloc()在全局存储器中分配空间两个参数：地址指针空间大小cudaFree()在全局存储器中回收空间一个参数回收空间地址指针,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设

36、计基于GPU的并行程序设计,第七章,2.6 CUDA设备存储器分配,示范代码分配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的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言2.CUDA编程模型2.1 CUDA编程接口

37、2.2 一个简单的CUDA程序2.3 线程的层次结构2.4 存储器的层次结构2.5 CUDA变量类型修饰符2.6 CUDA设备存储器分配2.7 CUDA Host-Device数据传输2.8 CUDA的函数声明3 CUDA并行编程示例,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.7 CUDA Host-Device 数据传输,cudaMemcpy()存储器数据传输参数：目的地址源地址传输字节数传输类型设备设备主机全局存储器全局存储器主机异步传输,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行

38、程序设计基于GPU的并行程序设计,第七章,2.7 CUDA Host-Device 数据传输,示范代码传输 64*64 单精度浮点数组M在主机存储器中Md在设备存储器中传输方向常数：cudaMemcpyHostToDevicecudaMemcpyDeviceToHost,cudaMemcpy(Md.elements,M.elements,size,cudaMemcpyHostToDevice);cudaMemcpy(M.elements,Md.elements,size,cudaMemcpyDeviceToHost);,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行

39、程序设计基于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的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.8 CUDA的函数声明,CUDA的函数声明_global_定义一个内核函数必须返回void_device_ 和 _host_可以一起使用,第二部

40、分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,2.8 CUDA的函数声明,CUDA的函数声明(续)_device_ 函数不能使用函数指针；_device_ 函数在设备上执行，所以:不能递归不能在函数内定义静态变量No variable number of arguments,调用一个内核函数线程生成,调用内核函数必须使用执行配置(execution configuration):,_global_ void KernelFunc(.);dim3 DimGrid(100,50);/5000 thread blocks dim3 Di

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

42、序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,3.1 通用编程策略,Global memory在设备内存中(DRAM)它的访问速度比shared memory要慢在设备上执行运算的优化方法是：把数据分割成子集（tile data），放入shared memory以一个线程块操作一个数据子集：从global memory读入数据子集到shared memory，使用多个线程实现存储级并行（memory-level parallelism）在shared memory中对数据子集执行运算;每一个线程可以有效率地多次访问任何数据元素从shared memor

43、y上拷贝结果到全局存储器,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,3.1 通用编程策略,Constant memory也在设备内存里(DRAM)比访问shared memory慢得多但是它可以高速缓存，对于只读数据的访问非常高效。通过访问模式小心划分数据R/Only constant memory(若在高速缓存内，非常快)R/W(块内共享)shared memory(非常快)R/W(线程内）registers(非常快)R/W（输入和结果）global memory(非常慢),第二部分 基于GPU的并行程序设计基于GPU

44、的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言2.CUDA编程模型3 CUDA并行编程示例3.1 通用编程策略3.2 CUDA编程框架3.3 矩阵相乘（使用全局内存）3.4 矩阵相乘（使用共享内存）,/全局变量声明_host_,_device_._constant_,_texture_/函数原型声明_global_ void kernelOne()/内核函数float handyFunction()/普通函数main()cudaMalloc(/从设备传输结果到主机_global_ void kernelOne(type args,)/内核函数/局部变量声

45、明_local_,_shared_/自动变量被默认分配到寄存器或本地存储器中float handyFunction(int inVar)/普通函数,CUDA编程框架,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言2.CUDA编程模型3 CUDA并行编程示例3.1 通用编程策略3.2 CUDA编程框架3.3 矩阵相乘（使用全局内存）3.4 矩阵相乘（使用共享内存）,代码参考：CUDA Programmng Guide（CUDA 2.0）,第二部分 基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计基于GPU的并行程序设计,第七章,内容,1.引言2.CUDA编程模型3 CUDA并行编程示例3.1 通用编程策略3.2 CUDA编程框架3.3 矩阵相乘（使用全局内存）3.4 矩阵相乘（使用共享内存）,代码参考：CUDA Programmng Guide（CUDA 2.0）,