Microsoft PowerPoint - 第xx章 GPU基础及CDUA编程简介

Size: px

Start display at page:

Download "Microsoft PowerPoint - 第xx章 GPU基础及CDUA编程简介"

勋出功
5 years ago
Views:

1 GPU 基础及 CUDA 编程简介王娴航天航空学院计算流体力学与实践内容 GPU 基础知识 CUDA 编程基础一维数组相加编程实例在计算流体力学中应用实例 --- 求解二维扩散方程矩阵相乘编程实例多 GPU 并行数据传递模式 2 1

2 1. GPU 基础知识什么是 GPU? GPU: Graphic Processing Unit 用于个人计算机工作站和游戏机的专用图像显示设备显示卡或主板集成 3 nvidia 和 ATI (now AMD) 是主要制造商 1. GPU 基础知识什么是 GPGPU? GPGPU: General Purpose Graphic Processing Unit 不用于图像处理显示, 一般用途总称 4 重要应用之一 : 工程上基于网格的计算流体力学 CFD 2

3 1. GPU 基础知识吸引力所在? 计算性能非常高 ALU: arithmetic logical unit 更多资源用于缓存更多资源用于数据计算 5 Core i7 : 4 cores 100GFlops TeslaK20 > 3.95TFlops 3950GFLOPS 2688 cores 1. GPU 基础知识 GPU-CPU 发展吸引力所在?

4 1. GPU 基础知识吸引力所在? 最新 GPU:Tesla K 系列 GPU 基础知识吸引力所在? 带宽非常高 250GB/s High memory bandwidth 8 4

5 1. GPU 基础知识吸引力所在? 9 1. GPU 基础知识吸引力所在? 价格低 ( 数千元 ~ 万元 ) 24.7cm TeslaK20: 约 2 万元性能 :3.95 TFLOPS 11.1cm 安装简单, 无需特殊环境,PC 安装 OK!! 10 桌面级超级计算机 5

6 1. GPU 基础知识怎么用? GPGPU 的综合开发环境 NVIDIA CUDA (Compute Unified device Architecture) 针对 NVIDIA 公司的 GPU 的编程软件, 可从该公司主页免费下载 ATI Stream DirectCompute OpenCL GPU 基础知识怎么用? Application Programming Interface 12 6

7 1. GPU 基础知识怎么用? GPU 基础知识怎么用? CUDA 源文件 a.cu 编译 : NVCC nvcc 只是个编译器驱动, 他会调用很多 cuda tools 去完成各个编译步骤, 最后调用系统的 c/c++ 编译器 (cl, gcc) 以及链接器 ( 所有非 CUDA 编译步骤扔给通用的 C 编译器 ) nvcc -o main.o -c -O3 a.cu -I${HOME}/NVIDIA_CUDA_SDK/common/inc gcc -o run main.o -O3 -L/usr/local/cuda/lib -lcudart -I 后面接的是源代码内包含的文件所在目录 (#include) -L 后面接的路径是函数库的搜索目录 14 -l 表示加入某个函数库 (cudart.lib) 7

1. GPU基础知识怎么用 CUDA是C语言扩展 CUDA Fortran是Fortran语言扩展 CUDA Fortran源文件 a.cuf 编译 PGI Fortran PGI CUDA Fortran Complier http://www.pgroup.com/resources/cudafortran.

8 1. GPU基础知识怎么用 CUDA是C语言扩展 CUDA Fortran是Fortran语言扩展 CUDA Fortran源文件 a.cuf 编译 PGI Fortran PGI CUDA Fortran Complier CUDA Fortran相关资源 CUDA Fortran Programming Guide and References 编程及API使用方法相关资料 15 产品 1. GPU基础知识怎么用 GPU本身不能工作由CPU控制 HPC CPU+GPU体系架构 GPU程序最大并行线程数 65535x65535x512 CPU 部分 Host code GPU 部分 Device code 并行 CPU 上执行程序串行部分 GPU上执行并行部分 16 8

9 1. GPU 基础知识 CPU- 串行线程管理? Grid,Block,Thread Maximum Thread: 65535x65535x512 GPU- 并行 CPU- 串行 GPU- 并行 GPU 基础知识 NVIDIA GPU 构造 8x30 = 240 (SP) 18 9

1. GPU 基础知识 NVIDIA GPU 构造 19 NVIDIA TeslaM2070 1.

chip CUDA program on CPU invokes kernel gird Streaming Processor Enumerating the

20 Video Memory SIMT:single-instruction, multiple thread 单指令多线程 Executing the

10 1. GPU 基础知识 NVIDIA GPU 构造 19 NVIDIA TeslaM GPU 基础知识 NVIDIA GPU 并行机理调用 CUDA kernel (GPU 程序 ) N 条线程 ( 同时 ) 执行该段程序 N 次 GPU chip CUDA program on CPU invokes kernel gird Streaming Processor Enumerating the blocks of a grid Streaming Multiprocessor Distributing the blocks of a grid to SM 20 Video Memory SIMT:single-instruction, multiple thread 单指令多线程 Executing the threads of a block concurrently on one MP SIMT MP maps each thread to one SP Each scalar thread executes independently 10

11 1. GPU 基础知识 GPU 基础知识 22 11

12 世界 HPC top 年 6 月中国天河二号中国天河一号中国 HPC top 10 12

13 1. GPU 基础知识世界十强全部是 CPU+GPU 体系架构 ; 国内排名前 3 的机器 2 套是 CPU+GPU 体系架构 ; 国内前 10 名里 4 台是 CPU+GPU 体系架构, 在 TOP100 中共有 13 套 CPU+GPU 异构机群 2. CUDA 编程基础 Maximum Limitation: Grid: Block: ( 512, 512, 64 ) Thread:

14 2. CUDA 编程基础并行线程组织结构 CUDA 编程基础并行线程组织结构内存冲突 28 14

15 2. CUDA 编程基础存储器模型寄存器共享内存全局内存 CUDA 编程基础内存分配 30 15

16 2. CUDA 编程基础 CPU-GPU 数据交换 CUDA 编程基础 CUDA 变量和函数 32 16

17 2. CUDA 编程基础 CUDA 变量和函数重要!! dimgrid.x = 2 dimgrid.y = 2 dimgrid.z = 1 dimblock.x = 4 dimblock.y = 2 dimblock.z = 2 共 2 x 2=4 个 Blocks 一个 Block: 4x 2 x 2=16 threads (MAX:512) 共 16x4=64 threads CUDA 编程基础 CUDA 变量和函数 34 17

18 2. CUDA 编程基础 CUDA 变量和函数 CUDA 编程基础 CUDA 变量和函数 36 18

2. CUDA 编程基础 CUDA 编程框架声明 GPU 核心程序 global global_ GPU( ); 在 CPU 和 GPU 上分配内存将数据从 CPU 内存复制至 GPU 显存用 dim3 指明 block 数目与 thread 数目调用 GPU 核心程序 Malloc( ), cudamallochost () & cudamalloc ( ) cudamemcpy(,

19 2. CUDA 编程基础 CUDA 编程框架声明 GPU 核心程序 global global_ GPU( ); 在 CPU 和 GPU 上分配内存将数据从 CPU 内存复制至 GPU 显存用 dim3 指明 block 数目与 thread 数目调用 GPU 核心程序 Malloc( ), cudamallochost () & cudamalloc ( ) cudamemcpy(, cudamemcpyhosttodevice) dim3 block (x,y,1), thread (m,n,p) GPU <<< block, thread >>> ( ) ::::::::: 将结果数据从 GPU 显存复制回 CPU 内存 cudamemcpy(, cudamemcpydevicetohost) 37 从 CPU 上输出结果 3. 一维数组相加 CUDA 编程实例 CPU A[1] A[2] A[3] A[n] B[1] B[2] B[3] B[n] = = = = C[1] C[2] C[3] C[n] 38 for (i=0; i<n; i++) C[i] = A[i] +B [i] ; 19

20 3. 一维数组相加 CUDA 编程实例 A[1] A[2] A[3] A[n] GPU B[1] B[2] B[3] B[n] = = = = C[1] C[2] C[3] C[n] 39 T1 T2 T3 Tn 3. 一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (256,1) ; dim3 dimgrid (??, 1); 一个元素对应一条线程 : 4 共 1024 条线程 GRID BLOCK BLOCK BLOCK BLOCK C(0) C(255) C(1023) 256 threads blockdim.x = threads blockdim.x = threads blockdim.x = threads blockdim.x = Blocks griddim.x =4 20

3. 一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (256,1) ; dim3 dimgrid ( 4, 1); 一个元素对应一条线程 : 共 1024 条线程 blockidx.

x = 0~255 threadidx.x = 0~255 threadidx.x = 0~255 i = blockdim.x x blockidx.x + threadidx.

一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (1,1) ; dim3 dimgrid ( 4, 1); 256 个元素对应一条线程 : 共 4 条线程 blockidx.x = 0 blockidx.

21 3. 一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (256,1) ; dim3 dimgrid ( 4, 1); 一个元素对应一条线程 : 共 1024 条线程 blockidx.x = 0 blockidx.x = 1 blockidx.x = 2 blockidx.x = threadidx.x = 0~255 threadidx.x = 0~255 threadidx.x = 0~255 threadidx.x = 0~255 i = blockdim.x x blockidx.x + threadidx.x ; C[i] = A[i] + B[i] ; 41 无循环 CPU 程序 : for (i=0; i<n; i++) C[i] = A[i] +B [i] ; 3. 一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (1,1) ; dim3 dimgrid ( 4, 1); 256 个元素对应一条线程 : 共 4 条线程 blockidx.x = 0 blockidx.x = 1 blockidx.x = 2 blockidx.x = threadidx.x = 0 threadidx.x = 0 threadidx.x = 0 threadidx.x = 0 42 for (k=0; k<256; k++) { i = 256 x blockidx.x + k ; C[i] = A[i] + B[i] ; } 每条线程计算 256 个元素 ( 即 256 次循环 ) 21

3. 一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (4,1) ; dim3 dimgrid ( 4, 1); blockidx.x = 0 256 个元素对应 4 条线程 : 共 16 条线程 43 0 1 2 3 0 1 2 3 0 1 2 3 threadidx.x = 0~3 blockidx.

22 3. 一维数组相加 CUDA 编程实例 n =1024 GPU dim3 dimblock (4,1) ; dim3 dimgrid ( 4, 1); blockidx.x = 个元素对应 4 条线程 : 共 16 条线程 threadidx.x = 0~3 blockidx.x = threadidx.x = 0~3 for (k=0; k<256; k++4) { i = 256 x blockidx.x + threadidx.x + k ; C[i] = A[i] + B[i] ; } 每条线程做 256/4 = 64 次循环 3. 一维数组相加 CUDA 编程实例 GPU 要点每个 Block 中的 Thread 数最大不超过 512; 为了充分利用 SM,Block 数尽可能多,>

23 3. 一维数组相加 CUDA 编程实例主程序 (cpu 上运行 ) 第一步 : 定义 Host 上矩阵 M N P 与相应的 Device 上矩阵 Md Nd Pd, 并将 Host 上矩阵数据拷贝至相应的 Device 矩阵 float *A,*Ad, *B,*Bd, *C,*Cd ; int size = n*sizeof(float); Malloc ((void**)&a, size ) ; ----CPU 上分配地址空间 cudamalloc ((void**)&ad, size ) ; ---GPU 上分配地址空间 45 cudamemcpy (Ad, A, size, cudamemcpyhosttodevice) 3. 一维数组相加 CUDA 编程实例 GPU 数组 B: Malloc ((void**)&b, size ) ; ----CPU 上分配地址空间 cudamalloc ((void**)&bd, size ) ; ---GPU 上分配地址空间 cudamemcpy (Ad, A, size, cudamemcpyhosttodevice) 数组 C: Malloc ((void**)&c, size ) ; ----CPU 上分配地址空间 cudamalloc ((void**)&cd, size ) ; ---GPU 上分配地址空间 46 ( 结果数组 Cd, 无需拷贝 C 中本无值 ) 23

24 3. 一维数组相加 CUDA 编程实例主程序 (cpu 上运行 ) 第二步 : 定义 block 大小, 即 Thread 数定义 Grid 大小, 即 Block 数 dim3 dimblock (256, 1) ; dim3 dimgrid (4, 1); 按照定义的 Thread Block 的大小启动 Device 程序 :Kernel 声明程序类型 : global void Plus(float*, float*, float*, int) ; 在 device 上执行, 仅由 host 调用 47 GPU 子程序调用 :Kernel Plus <<<dimgrid, dimblock >>> (Ad, Bd, Cd, n) ; 3. 一维数组相加 CUDA 编程实例主程序 (cpu 上运行 ) 第三步 : 将计算结果从 Device 拷贝至 Host 释放 Device 内存空间 cudamemcpy (C, Cd, size, cudamemcpydevicetohost) free(a); free(b); free(c); cudafree(ad); cudafree(bd); cudafree(cd); 48 24

25 global void Plus (float*, float*, float*, int) ; 49 int main(int argc, char *argv[]) { float *A,*Ad, *B,*Bd, *C,*Cd; int size = n*sizeof(float); Malloc ((void**)&a, size ) ; Malloc ((void**)&b, size ) ; Malloc ((void**)&c, size ) ; cudamalloc ((void**)&ad, size ) ; cudamalloc ((void**)&bd, size ) ; cudamalloc ((void**)&cd, size ) ; cudamemcpy (Ad, A, size, cudamemcpyhosttodevice) ; cudamemcpy (Bd, B, size, cudamemcpyhosttodevice) ; cudamemcpy (Cd, C, size, cudamemcpyhosttodevice) ; dim3 dimblock (256, 1) ; dim3 dimgrid (4, 1); Plus <<<dimgrid, dimblock >>> (Ad, Bd, Cd, n) ; cudamemcpy (C, Cd, size, cudamemcpydevicetohost); 主程序 free(a); free(b); free(c); cudafree(ad); cudafree(bd); cudafree(cd); } 3. 一维数组相加 CUDA 编程实例 global void Plus (float* A, float* B, float* C, int n) { } int i; i = blockdim.x x blockidx.x + threadidx.x ; C[i] = A[i] + B[i] ; GPU 子程序 50 25

26 注意 3. 一维数组相加 CUDA 编程实例 dimblock,dimgrid 是 dim3 型变量, 有两个分量 : dim3 dimblock (256,1) ; dimblock.x =256 dimblock.y =1 dim3 dimgrid ( 4, 1); dimgrid.x =4 dimgrid.y=1 可以改成任何名字 BLOCK 如 :dim3 aa(256,1) ; dim3 bb(4,1) C(0) C(255) aa.x=256, aa.y=1; bb.x=4, bb.y= threads blockdim.x = 256 blockidx.x = threadidx.x = 0~255 Plus <<<dimgrid, dimblock >>> (Ad, Bd, Cd, n) ; Plus <<<aa, bb >>> (Ad, Bd, Cd, n) ; griddim.x,blockdim.x,blockidx.x,threadidx.x ( 或.y 或.z) 等是 GPU 根据 kernel 中的 dim 变量自行确定的名字不能更改子程序中只能出现这些变量, 而不能出现 dim3 型变量 dimblock.x, aa.x 等. 4. 在 CFD 中应用实例二维扩散方程 52 26

27 4. 在 CFD 中应用实例二维扩散方程 A typical stencil computing 在 CFD 中应用实例二维扩散方程 54 27

28 4. 在 CFD 中应用实例 CPU 程序 % 求余数 ( 周期边界 ) 在 CFD 中应用实例 CUDA 程序 GPU 上只能分配 1 维数组 CPU 侧 GPU 侧计算变量 F, Fn f fn 56 28

在 CFD 中应用实例 CUDA 程序 ny [jx, jy] 58 f cpu nx [ jy][ jx] f [ nx jy jx] gpu

29 4. 在 CFD 中应用实例 CUDA 程序对比 CPU 程序声明 :mykernelfunction_cpu (float *x, float *y, float alpha) 对比调用 CPU 子程序 :mykernelfunction_cpu (x,y,alpha) 在 CFD 中应用实例 CUDA 程序 ny [jx, jy] 58 f cpu nx [ jy][ jx] f [ nx jy jx] gpu jy blockdimy. blockidxy. threadidxy. jx blockdim. x blockidx. x threadidxx. 29

31 nx=

32 ny=256 nx=

33 65 4. 在 CFD 中应用实例 CUDA 程序 Data transfer from Device (GPU) to Host (CPU) cudamemcpy (F, f, sizeof(float)*n, cudamemcpydevicetohost); cudamemcpy (Fn, fn, sizeof(float)*n, cudamemcpydevicetohost); CPU 侧 GPU 侧 CPU Output results F, Fn 计算变量 F, Fn f fn 66 33

34 4. 在 CFD 中应用实例 CUDA 程序矩阵相乘 CUDA 编程实例 CUDA 程序中线程和内存的管理共享内存的使用 shared memory 寄存器的使用 register 线程 ID Host 和 Device 之间数据传递 68 34

35 5. 矩阵相乘 CUDA 编程实例矩阵 : M: 宽度 ( 列 )wm, 高度 ( 行 )hm N wn hn N: 宽度 ( 列 )wn, 高度 ( 行 )hn M P P = M x N hm hm P: 宽度 ( 列 )wn, 高度 ( 行 )hm wm wn 矩阵相乘 CUDA 编程实例 N 方阵 M N:WIDTH( 各方向上元素个数 ) P = M x N 方阵 WIDTH WIDTH 一条线程处理 P 的一个单元总线程数 : M P WIDTH xwidth WIDTH WIDTH 70 WIDTH 35

36 5. 矩阵相乘 CUDA 编程实例主程序第一步 : 定义 Host 上矩阵 M N P 与相应的 Device 上矩阵 Md Nd Pd, 并将 Host 上矩阵数据拷贝至相应的 Device 矩阵矩阵相乘 CUDA 编程实例矩阵 M 与 Md CPU 上运行程序部分 float *M,*Md ; int size = WIDTH*WIDTH*sizeof(float); Malloc ((void**)&m, size ) ; ----CPU 上分配地址空间 cudamalloc ((void**)&md, size ) ; ---GPU 上分配地址空间 cudamemcpy (Md, M, size, cudamemcpyhosttodevice) 72 36

37 5. 矩阵相乘 CUDA 编程实例矩阵 N 与 Nd CPU 上运行程序部分 float *N,*Nd ; int size = WIDTH*WIDTH*sizeof(float); Malloc ((void**)&n, size ) ; ----CPU 上分配地址空间 cudamalloc ((void**)&nd, size ) ; ---GPU 上分配地址空间 cudamemcpy (Nd, N, size, cudamemcpyhosttodevice) 矩阵相乘 CUDA 编程实例结果矩阵 P 与 Pd CPU 上运行程序部分 float *P,*Pd ; int size = WIDTH*WIDTH*sizeof(float); Malloc ((void**)&p, size ) ; ----CPU 上分配地址空间 cudamalloc ((void**)&pd, size ) ; ---GPU 上分配地址空间 cudamemcpy (Pd, P, size, cudamemcpyhosttodevice) 74 37

38 5. 矩阵相乘 CUDA 编程实例主程序第二步 : 定义 block 大小, 即 Thread 数定义 Grid 大小, 即 Block 数按照定义的 Thread Block 的大小启动 Device 程序 :Kernel 矩阵相乘 CUDA 编程实例 CPU 上运行程序部分 dim3 dimblock (BLOCK_SIZE, BLOCK_SIZE) ; dim3 dimgrid (WIDTH/dimBlock.x, WIDTH/dimBlock.y); x 方向的 Thread 数 :dimblock.x = BLOCK_SIZE y 方向的 Thread 数 :dimblock.y = BLOCK_SIZE 一个 Block 中的总 Thread 数 : BLOCK_SIZE x BLOCK_SIZE 76 x 方向的 Block 数 :dimgrid.x = WIDTH/dimBlock.x y 方向的 Block 数 :dimgrid.y = WIDTH/dimBlock.x 总 Block 数 :WIDTH/dimBlock.x x WIDTH/dimBlock.x 38

39 5. 矩阵相乘 CUDA 编程实例 CPU 上运行程序部分声明程序类型 : global void Muld(float*, float*, float*, int) ; 在 device 上执行, 仅由 host 调用 GPU 子程序调用 :Kernel Muld <<<dimgrid, dimblock >>> (Md, Nd, Pd, WIDTH) ; 矩阵相乘 CUDA 编程实例主程序第三步 : 将计算结果从 Device 拷贝至 Host 释放 Device 内存空间 cudamemcpy (P, Pd, size, cudamemcpydevicetohost) free(m); free(n); free(p); 78 cudafree(md); cudafree(nd); cudafree(pd); 39

40 79 #define BLOCK_SIZE 16 global void Muld (float*, float*, float*, int) ; int main(int argc, char *argv[]) { float *M,*Md, *N,*Nd, *P,*Pd; int size = WIDTH*WIDTH*sizeof(float); Malloc ((void**)&m, size ) ; Malloc ((void**)&n, size ) ; Malloc ((void**)&p, size ) ; cudamalloc ((void**)&md, size ) ; cudamalloc ((void**)&nd, size ) ; cudamalloc ((void**)&pd, size ) ; cudamemcpy (Md, M, size, cudamemcpyhosttodevice) ; cudamemcpy (Nd, N, size, cudamemcpyhosttodevice) ; cudamemcpy (Pd, P, size, cudamemcpyhosttodevice) ; dim3 dimblock (BLOCK_SIZE, BLOCK_SIZE) ; dim3 dimgrid (WIDTH/dimBlock.x, WIDTH/dimBlock.y); Muld <<<dimgrid, dimblock >>> (Md, Nd, Pd, WIDTH) ; cudamemcpy (P, Pd, size, cudamemcpydevicetohost); 主程序 free(m); free(n); free(p); cudafree(md); cudafree(nd); cudafree(pd); } 矩阵相乘 CUDA 编程实例 N 求解的 Block 编号 ( 自动生成 ): bx = blockidx.x by = blockidx.y Blocks 在 MP 上并列计算一条 Thread 求解 P 矩阵的一个元素,Thread 编号 ( 自动生成 ): tx = threadidx.x ty = threadidx.y M P Threads 在 SP 上并列计算本图示 : bx=2,by=2 tx=1,ty=1 ( 从 0 开始 ) 80 Block 40

41 矩阵相乘 CUDA 编程实例 N 子程序思路外循环 :Block, 步长 BLOCK_SIZE 内循环 :Thread, 步长 1,k=k+1 BLOCK_SIZE 行 k M BLOCK_SIZE 列 P k 81 Block 矩阵相乘 CUDA 编程实例外循环起点终点及步长 N abegin = WIDTH * BLOCK_SIZE * by ; aend = abegin+width-1 ; bstep bbegin = BLOCK_SIZE * bx ; bx 块 astep = BLOCK_SIZE ; bstep 之间相差元素 bstep = BLOCK_SIZE * WIDTH; BLOCK_SIZE 个元素 bbegin n 列与 n+k 列之间元素编号相差 M k P n 行与 n+k 行之间元素编号相差 astep (k*width) abegin by 块 aend 82 BLOCK_SIZE 个元素 WIDTH 个元素 Block 41

42 实现外循环共享内存 N for (a = abegin, b = bbegin; a<=aend; a+=astep, b+=bstep) a,b 变, 即 Block 变 { shared float Ms [BLOCK_SIZE] [BLOCK_SIZE] ; shared float Ns [BLOCK_SIZE] [BLOCK_SIZE] ; Ms[ty][tx] =M[a + WIDTH *ty + tx] ; Ns[ty][tx] =N [b + WIDTH *ty + tx] ; b 外循环 :Block 步长 : BLOCK_SIZE 注意 : Ms,Ns 是二维坐标 M,N 一维坐标 a,b 为总体一维坐 83 标的编号 M a P Block 实现内循环 N 内循环 :Thread, 步长 1,k=k+1 for (int k=0; k<block_size; k++) Psub += Ms[ty][k] * Ns[k][tx] ; k 外循环的乘积相加的值也累积于 Psub 上 M k P 84 Block 42

43 5. 矩阵相乘 CUDA 编程实例将在共享内存中计算好的元素写回显存注意 : 一维与二维数组的对应关系! int c = WIDTH * BLOCK_SIZE * by + BLOCK_SIZE *bx; P[ c + WIDTH*ty + tx] = Psub; P 所求元素的总体一维编号 BLOCK_SIZE bx c by 85 BLOCK_SIZE WIDTH global void Muld (float* M, float* N, float* P, intwidth) { float Psub; int a,b,c,k; bx = blockidx.x; by = blockidx.y; tx = threadidx.x; ty = threadidx.y; abegin = WIDTH * BLOCK_SIZE * by ; aend = abegin+width-1 ; bbegin = BLOCK_SIZE * bx ; astep = BLOCK_SIZE ; bstep = BLOCK_SIZE * WIDTH; GPU 子程序 86 } for (a = abegin, b = bbegin; a<=aend; a+=astep, b+=bstep) { shared float Ms [BLOCK_SIZE] [BLOCK_SIZE] ; shared float Ns [BLOCK_SIZE] [BLOCK_SIZE] ; } Ms[ty][tx] =M[a + WIDTH *ty + tx] ; Ns[ty][tx] =N [b + WIDTH *ty + tx] ; syncthreads() ; for (int k=0; k<block_size; k++) Psub += Ms[ty][k] + Ns[k][tx] ; syncthreads() ; c = WIDTH * BLOCK_SIZE * by + BLOCK_SIZE *bx; P[ c + WIDTH*ty + tx] = Psub; 43

44 5. 矩阵相乘 CUDA 编程实例使用共享内存 N 中 BLOCK 一条 Thread 对应 P 的一个元素一个 Block 中的所有线程对共享内存均可视如 : 求解 T1 T8 均用到 M 矩阵中 a1 a8 的值不用共享内存 : 需要以速率 v 从 Global Memory 中将 a1-a8 调用 8 次 ; 使用共享内存 : 仅需要以速率 v 调用 1 次, 其余均以约 100v 从共享内存中调用 M 中 BLOCK a1 a2 a8 P 中 BLOCK T1 T2 T 矩阵相乘 CUDA 编程实例使用共享内存一条 Thread 对应 P 的一个元素一个 Block 中的所有线程对共享内存均可视 T1 T5 Chip T1 约 GB/s T5 Chip T2 T3 shared T6 T7 T2 T3 shared a1-a8 T6 T7 T4 T8 T4 T GB/s a1-a8 Global Memory 100 GB/s a1-a8 Global Memory 44

5. 多 GPU 并行数据传递模式 GPU 间无法自行数据通讯 ; GPU 之间的数据交换需通过 CPU 来完成

cudamemcpydevicetohost X 89 cudamemcpyhosttodevice 5.

cudamemcpydevicetohost GPU MPI_Send CPU CPU MPI_Recv

45 5. 多 GPU 并行数据传递模式 GPU 间无法自行数据通讯 ; GPU 之间的数据交换需通过 CPU 来完成 cudamemcpydevicetohost MPI_Send MPI_Recv cudamemcpyhosttodevice cudamemcpydevicetohost X 89 cudamemcpyhosttodevice 5. 多 GPU 并行数据传递模式 HPC 模式 :CPU 数 >=GPU 数 CPU cudamemcpyhosttodevice cudamemcpydevicetohost GPU MPI_Send CPU CPU MPI_Recv cudamemcpyhosttodevice cudamemcpydevicetohost MPI_Send MPI_Recv cudamemcpyhosttodevice cudamemcpydevicetohost GPU GPU

46 5. 多 GPU 并行数据传递模式 HPC 模式 :CPU 数 >=GPU 数 node1 node2 node cudasetdevice(); myid =0 myid =1 myid =2 myid =3 myid =4 myid =5 二维求解区域 MPI 并行 MPI_Comm_rank(MPI_COMM_WORLD,&myid); 自动生成 myid 91 1GPU/node, 共用 6 个 nodes: cudasetdevice(0); 2GPU/node, 共用 3 个 nodes: cudasetdevice(myid%2); 3GPU/node, 共用 2 个 nodes: cudasetdevice(myid%3); 46

NSCC-TJ汇报

NSCC-TJ汇报 GPU 与 CUDA 简介赵洋工程师 zhaoyang@nscc-tj.gov.cn 国家超级计算天津中心内容提要天河一号架构 GPU 简介 CUDA 硬件模型和编程模型 CUDA Fortran & CUDA C CUDA 程序优化策略天河一号超级计算机系统 7168 computing nodes 1024 service nodes monitoring and diagnostic