并行程序设计Chapter-6.ppt

1、第六章 异构加速部件的并行编程,刘 轶 北京航空航天大学 计算机学院,6.1 Intel MIC协处理器编程简介MIC-Many Integrated Core商品名：Intel Xeon Phi,6.1 Intel MIC协处理器编程简介,一、Intel MIC硬件架构 第一代MIC: Knights Corner 60核，支持240个并行线程 共享内存，核间cache一致性 单处理器性能达1TeraFlops双精度浮点,第一代MIC核的特性 基于Pentium架构，做了改进(64位，SIMD支持) 32KB L1 cache, 512KB L2 cache(核内) 支持512bit向量，传

2、统的SSE 128bit，AVX 256bit 支持4个超线程，普通Xeon处理器是2个超线程 主频1.053GHz,为Host system + Coprocessor卡形式 8GB内存，320GB/s访存带宽 Linux系统，IP addressable 功耗225W Host上运行完整的Linux系统 Xeon Phi上运行一个嵌入式Linux，以及与host通信的驱动,新一代MIC: Knights Landing 72核 单精度浮点性能：6TeraFlops，双精度浮点性能：3TeraFlops 可独立工作,二、软件及编程语言 可为其配置IP地址 在host系统上可ssh到其上 编程

3、语言与接口基本与Intel Xeon处理器一致 编程语言 C/C+、Fortran等 编程接口 OpenMP Cilk TBB pthreads ,在ssh窗口中，显示/proc/cpuinfo文件,三、Intel MIC的几种计算模型,Offload模式 Symmetric模式 Native模式,6.1 Intel MIC协处理器编程简介,四、基本语法 最基本的关键字：offload #pragma offload 紧跟该命令之后的代码块(单条语句或)将在MIC上执行 offload只起到将代码移至MIC执行的作用，要使用多核并行执行，需配合使用OpenMP等并行编程接口/语言,在MIC的一

4、个核上执行for循环,使用OpenMP生成多线程，在MIC的多个核上并行执行for循环,6.1 Intel MIC协处理器编程简介,四、基本语法 数据传输关键字 in: 在设备端开辟空间，将主机端数据复制到设备端 out: 在设备端开辟空间，设备端代码执行结束后，将数据复制到主机端 inout: 在设备端开辟空间，执行设备端代码时，将数据复制到设备端，设备端代码执行完毕后，将数据复制到主机端 nocopy: 仅建立空间，不复制数据,在MIC计算出数组attr，需要将结果输出到主机端,MIC协处理卡上有独立内存，不与主机端共享内存，因此需进行主机MIC和MIC主机间的数据传输,注：非数组变量可不

5、显式地传输，自动被认为是inout,6.1 Intel MIC协处理器编程简介,四、基本语法 数据传输关键字：语法规则 传输关键字后跟( )，括号内为一个或多个变量名，之间逗号隔开 变量名应为数组、指针或普通变量 变量为数组或指向数组的指针时，可指定起始下标和长度 变量为指针时，只能指向非指针变量，且后跟length关键字指定元素个数 例：in(c:length(20)，或in(a,b,c:length(20) 元素个数也可以是变量，如in(c:lengthlen) 变量后除length关键字外，还可使用以下关键字(与length一样前加:) alloc_if和free_if：参数为表达式，如

6、结果为真，则进入/离开设备端时为变量开辟/释放空间 align：参数为正整数，应为2、4、8、16、，指定设备端开辟前述变量的对齐长度 alloc：参数为数组名，创建指定的部分内存空间 into：参数为变量或数组名，只能一对一传递，用于将数组从主机端拷贝到设备端的另一个数组，或相反,6.1 Intel MIC协处理器编程简介,代码示例一,数组a是MIC上使用的中间变量 使用nocopy关键字可避免在主机端和设备端之间进行数据传输 分配空间后不释放，在后续MIC代码中可继续使用,6.1 Intel MIC协处理器编程简介,代码示例二,使用in/out/inout可只传输数组的一部分数据，减少数据

7、传输时间和MIC端内存占用 MIC将开辟数组从第1个元素开始的内存空间，例如in(p10:100)将在MIC端为p0p109开辟空间，但不传输p0p9,6.1 Intel MIC协处理器编程简介,四、基本语法 关于taget关键字 使用target(mic)时，如MIC卡状态不正常(如启动失败)，将在CPU上执行代码，如程序员未察觉，将可能产生误导 使用target(mic)且系统中有多块MIC时，系统将轮流使用，各块卡的初始化时间都将被计入执行时间，可能导致程序执行时间变长,6.1 Intel MIC协处理器编程简介,四、基本语法 关于taget关键字 假设系统中有2块MIC卡，以上程序将创

8、建3个线程，线程0和线程1分别在两块MIC卡上执行，线程0在CPU上执行,6.1 Intel MIC协处理器编程简介,四、基本语法 异步执行与异步传输 可使用signal/wait关键字实现MIC和CPU代码的异步执行,异步执行与异步传输 使用offload_transfer关键字数据的异步传输,6.2 面向GPU的CUDA编程简介,6.2 面向GPU的CUDA编程简介,一、GPU硬件架构 Nvidia Tesla K20X主要参数 2688个流处理器核 处理器频率：732MHz 单精度浮点性能：3.95TFlops 双精度浮点性能：1.31TFlops 专用显存数量：6GB 最大功耗：235

9、Watt 接口：PCI Express x16,Nvidia Tesla K20X组成框图 (核心芯片：GK110),6.2 面向GPU的CUDA编程简介,一、GPU硬件架构 GK110主要构成 15个SMX(Stream Multiprocessor) 依型号不同可能有13或14个 K20X: 14 x 192 = 2688核 每个SMX 包含192个核 任务调度以32个线程为一组，称为warp，可同时调度4个warp执行 64KB片上内存 可配置为48KB shared memory+16KB L1 cache或16KB L1 cache + 48KB shared memory,SMX内

10、存层次结构,GK110由15个SMX构成,6.2 面向GPU的CUDA编程简介,与编程相关的主要参数指标,6.2 面向GPU的CUDA编程简介,二、软件及编程语言 编程语言：C、Fortran 编程接口：CUDA、OpenCL、DirectCompute,Nvidia GPU的编程语言/接口,6.2 面向GPU的CUDA编程简介,三、简单的程序 示例程序一：Hello world,_global_Cuda C的修饰符，表示该函数在GPU上运行，编译器需将其编译为设备代码(device code) kernel()在设备端执行kernel函数，尖括号中参数为传递给设备运行时环境的参数,注：核函数

11、在GPU上串行执行,示例程序二：整数加,cudaMalloc()和cudaFree()Cuda C接口函数，类似于标准C的malloc()和free()函数，用于在设备上分配和释放内存 cudaMemcpy()Cuda C接口函数，类似于标准C的memcpy()，但增加了一个参数用于指定源/目的指针分别是主机/设备指针,注意：对设备端分配的指针(即CudaMalloc()分配的指针，本例中为dev_c)： 主机端代码不得直接读/写该指针所指向内容，须用CudaMemcpy()拷贝回主存后才可读/写 主机端代码可将其作为参数传递，也可对指针进行运算,四、GPU并行程序 示例程序三：向量求和,ad

12、d()N: 执行核函数时使用的并行线程块(thread block)的数量1: 每个线程块中的线程数量 执行该核函数线程个数 = N1 blockIdx.xCuda C预定义数据结构，其中存放线程块的序号，在本例中，其取值范围为0, 1, , N-1 在本例中，核函数add启动了N个线程块，每个线程块包含1个线程，每个线程完成一个向量元素的加法,6.2 面向GPU的CUDA编程简介,五、并行的层次结构 Thread: 并行的基本单位 GPU线程要比CPU线程轻量得多 Thread block: 互相合作的线程组 Cooperative Thread Array (CTA) 允许彼此同步 通过快

13、速共享内存交换数据 以1维、2维或3维组织 通常最多包含512或1024个线程 Grid: 一组thread block 以1维、2维或3维组织 共享全局内存 Kernel: 在GPU上执行的核函数 一个核函数对应一个Grid,五、GPU并行程序 向量求和程序的另一种线程划分方案,与前一个程序版本的区别启动1个线程块，其中包含N个线程 threadIdx.xCuda C预定义数据结构，其中存放线程的序号，在本例中，其取值范围为0, 1, , N-1,同一线程块内的线程 vs. 不同线程块的线程，有何区别？,6.2 面向GPU的CUDA编程简介,五、并行的层次结构 运行环境为每个线程块创建独立的

14、副本，包括变量 不同线程块不直接共享变量 一个线程块内的多个线程使用同一个副本 同一线程块内的线程可直接共享变量和相互协作 一个线程块内的线程个数有上限值 上限值定义在设备属性结构中的maxThreadsPerBlock字段中 通常为512或1024，即一个线程块最多拥有512或1024个线程 threadId可以按一维、二维组织，blockID可以按一维、二维、三维组织 便于对多维数据(如图像或矩阵)进行并行处理,以前面的向量求和程序为例，如果向量长度超过512，就需要使用多个线程块进行计算： 核函数中线程索引的计算：tid = threadIdx.x + blockIdx.x * bloc

15、kDim.x; 核函数的调用(假设每个线程块中线程数量固定为128个)：add( dev_a, dev_b, dev_c );,6.2 面向GPU的CUDA编程简介,六、GPU内存结构及使用 线程可以: 读/写本线程寄存器(Registers) 读/写本线程Local memory 读/写本线程块Shared memory 读/写本Grid的Global memory 读本Grid的Constant memory (可缓存) 读本Grid的Texture memory (可缓存) cudaMalloc/cudaFree用于分配/释放Global memory cudaMemcpy在Host和D

16、evice间传输数据 常量内存(Constant memory)和纹理内存(Texture memory)中的内容在程序执行期间不变 可被加载到cache中，访问性能优于global memory,(Device) Grid,Constant Memory,Texture Memory,Global Memory,Block (0, 0),Shared Memory,Local Memory,Thread (0, 0),Registers,Local Memory,Thread (1, 0),Registers,Block (1, 0),Shared Memory,Local Memory,T

17、hread (0, 0),Registers,Local Memory,Thread (1, 0),Registers,Host,cudaMemcpy(),访问延迟=1cycle,访问延迟500cycle,六、GPU内存结构及使用 用于变量的描述符: _device_ 储存于global memory 生命期(lifetime)与应用程序相同 可被grid中所有线程存取, CPU代码通过runtime函数存取 _constant_ 储存于constant memory 生命期(lifetime)与应用程序相同 可被grid中所有线程读取, CPU代码通过runtime函数存取 _shared_ 储存于thread block的shared memory 生命期(lifetime)与thread block相同 只能被thread block内的线程存取 无修饰(Local变量) 储存于SM内的寄存器和local memory 生命期(lifetime)与线程相同 Thread私有,(De