第6章

Size: px

Start display at page:

Download "第6章"

牙井
4 years ago
Views:

1 计算机体系结构周学海中国科学技术大学

2 05/15-Review GPU: 多线程协处理器 GPU 编程模型 :SPMD (Single Program Multiple Data) 使用线程 (SPMD 编程模型 ), 不是用 SIMD 指令编程每个线程执行同样的代码, 但操作不同的数据元素每个线程有自己的上下文 ( 即可以独立地启动 / 执行等 ) 计算由大量的相互独立的线程 (CUDA threads or microthreads) 完成, 这些线程组合成线程块 (thread blocks) GPU 执行模型 :SIMT (Single Instruction Multiple Thread) 一组执行相同指令的线程由硬件动态组织成 warp 一个 warp 是由硬件形成的 SIMD 操作 GPU 存储器组织 Local Memory, Shared Memory, Global Memory GPU 分支处理 ( 发散与汇聚 ) 5/15/2019 中国科学技术大学 2

3 5/15/2019 中国科学技术大学 3

4 Traditional CPU code 5/15/2019 中国科学技术大学 4

5 CUDA code 5/15/2019 中国科学技术大学 5

6 GPU: a multithreaded coprocessor SP: scalar processor CUDA core Executes one thread SM streaming multiprocessor 32xSP (or 16, 48 or more) Fast local shared memory (shared between SPs) 16 KiB (or 64 KiB) SM SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SP SHARED MEMORY GLOBAL MEMORY (ON DEVICE) 5/15/2019 中国科学技术大学 6

7 Warps are multithreaded on core 一个 warp 由 32 个 µthreads 构成多个 warp 线程在单个核上交叉运行, 以隐藏存储器访问和功能部件的延迟单个线程块包含多个 warp, 这些 warp 都映射到同一个核上多个线程块也可以在同一个核上运行 5/15/2019 [Nvidia, 2010] 中国科学技术大学 7

8 5/15/2019 中国科学技术大学

9 Threads Can Take Different Paths in Warp-based SIMD 每个线程可以包含控制流指令这些线程可以执行不同的控制流路径 A C B D F Thread Warp Thread 1 Thread 2 Thread 3 Common PC Thread 4 E G 5/15/2019 中国科学技术大学 9 Slide credit: Tor Aamodt

10 Control Flow Problem in GPUs/SIMT GPU 控制逻辑使用 SIMD 流水线以节省资源这些标量线程构成 warp 当一个 WARP 中的线程分支到不同的执行路径时, 产生分支发散 (Branch divergence) Branch Path A Path B 与向量处理机模型的条件执行类似 (Vector Mask and Masked Vector Operations?) 5/15/2019 中国科学技术大学 10

11 Conditional Branching 与向量结构类似, GPU 使用内部的屏蔽字 (masks) 还使用了 (SIMT-Stack) 分支同步堆栈保存分支的路径地址保存该路径的 SIMD lane 屏蔽字 (mask) 即指示哪些车道可以提交结果指令标记 (instruction markers) 管理何时分支 (divergence) 到多个执行路径, 何时路径汇合 (converge) PTX 层 CUDA 线程的控制流由 PTX 分支指令 (branch call return and exit) 控制以及由程序员指定的每个线程车道的 1-bit 谓词寄存器 GPU 硬件指令层, 控制流包括 : 分支指令 (branch,jump call return) 特殊的指令用于管理分支同步栈 GPU 硬件为每个 SIMD thread 提供堆栈保存分支的路径 GPU 硬件指令带有控制每个线程车道的 1-bit 谓词寄存器 5/15/2019 中国科学技术大学

12 Branch divergence 硬件跟踪各 µthreads 转移的方向 ( 判定哪些是成功的转移, 哪些是失败的转移 ) 如果所有线程所走的路径相同, 那么可以保持这种 SIMD 执行模式如果各线程选择的方向不一致, 那么创建一个屏蔽 (mask) 向量来指示各线程的转移方向 ( 成功失败 ) 继续执行分支失败的路径, 将分支成功的路径压入硬件堆栈 ( 分支同步堆栈 ), 待后续执行 SIMD 车道何时执行分支同步堆栈中的路径? 通过执行 pop 操作, 弹出执行路径以及屏蔽字, 执行该转移路径 SIMD lane 完成整个分支路径执行后再执行下一条指令称为 converge( 汇聚 ) 对于相同长度的路径,IF-THEN-ELSE 操作的效率平均为 50% 5/15/2019 中国科学技术大学

13 Example if (X[i]!= 0) X[i] = X[i] Y[i]; else X[i] = Z[i]; ld.global.f64 RD0, [X+R8] ; RD0 = X[i] setp.neq.s32 P1, RD0, #0 ; P1 is predicate register bra ELSE1, *Push ; Push old mask, set new mask bits ; if P1 false, go to ELSE1 ld.global.f64 RD2, [Y+R8] ; RD2 = Y[i] sub.f64 RD0, RD0, RD2 ; Difference in RD0 st.global.f64 [X+R8], RD0 ; X[i] = bra ENDIF1, *Comp ; complement mask bits ; if P1 true, go to ENDIF1 ELSE1: ld.global.f64 RD0, [Z+R8] ; RD0 = Z[i] st.global.f64 [X+R8], RD0 ; X[i] = RD0 ENDIF1: <next instruction>, *Pop ; pop to restore old mask 5/15/2019 中国科学技术大学 13

14 Branch Divergence Handling A; if (some condition) { B; } else { C; } D; A TOS One per warp Control Flow Stack Next PC Recv PC Active Mask DA B D 1110 C D 0001 Execution Sequence B D C A C B D Time 5/15/2019 中国科学技术大学 14

15 SIMT stack operation Example CUDA C source code for illustrating SIMT stack operation Example PTX assembly code for illustrating SIMT stack operation. 5/15/2019 中国科学技术大学 15

16 5/15/2019 中国科学技术大学 16

17 Remember: Each Thread Is Independent SIMT 主要优点 : 可以独立地处理线程, 即每个线程可以在任何标量流水线上单独执行 ( MIMD 处理模式 ) 可以将线程组织成 warp, 即可以将执行相同指令流的线程构成 warp, 形成 SIMD 处理模式, 以充分发挥 SIMD 处理的优势如果有许多线程, 对具有相同 PC 值的线程可以将它们动态组织到一个 warp 中这样可以减少分支发散提高 SIMD 利用率 SIMD 利用率 : 执行有用操作的 SIMD lanes 的比例 ( 即, 执行活动线程的比例 ) 5/15/2019 中国科学技术大学 17

18 Dynamic Warp Formation/Merging Idea: 分支发散之后, 动态合并执行相同指令的线程从那些等待的 warp 中形成新的 warp 足够多的线程分支到每个路径, 有可能创建完整的新 warp Warp X Warp Y Warp Z 5/15/2019 中国科学技术大学 18

19 Dynamic Warp Formation/Merging Idea: 分支发散之后, 动态合并执行相同指令的线程 Branch Path A Path B Fung et al., Dynamic Warp Formation and Scheduling for Efficient GPU Control Flow, MICRO /15/2019 中国科学技术大学 19

20 Dynamic Warp Formation Example B x/1110 y/0011 A x/1111 y/1111 C x/1000 y/0010 D x/0110 y/0001 F x/0001 y/1100 E x/1110 y/0011 A D Execution of Warp x at Basic Block A Legend A A new warp created from scalar threads of both Warp x and y executing at Basic Block D Execution of Warp y at Basic Block A Baseline Dynamic Warp Formation G x/1111 y/1111 A A B B C C D D E E F F G G A A A A B B C D E E F G G A A Time Time 5/15/2019 中国科学技术大学 20

21 Hardware Constraints Limit Flexibility of Warp Grouping Functional Unit Registers for each Thread Registers for thread IDs 0, 4, 8, Registers for thread IDs 1, 5, 9, Registers for thread IDs 2, 6, 10, Registers for thread IDs 3, 7, 11, Lane Can you move any thread flexibly to any lane? Memory Subsystem 5/15/2019 中国科学技术大学 21

22 When You Group Threads Dynamically 存储器访问如何处理? 固定模式的存储器访问相对简单, 当动态构成 warp 时, 使得访问模式具有随机性, 使得问题变得复杂降低存储器访问的局部性导致存储器带宽利用率的下降 5/15/2019 中国科学技术大学 22

23 What About Memory Divergence? 现代 GPUs 包括高速缓存, 减少对存储器的访问 Ideally: 一个 warp 中的所有线程的存储器访问都命中 ( 互相没有冲突 ) Problem: 一个 Warp 中有些命中, 有些失效 Problem: 一个线程的 stall 导致整个 warp 停顿需要有相关技术来解决存储器发散访问问题 5/15/2019 中国科学技术大学 23

24 NVIDIA GeForce GTX 285 NVIDIA-speak: 240 stream processors SIMT execution Generic speak: 30 cores 8 SIMD functional units per core 5/15/2019 中国科学技术大学 24

25 NVIDIA GeForce GTX 285 core 64 KB of storage for thread contexts (registers) = SIMD functional unit, control shared across 8 units = multiply-add = multiply = instruction stream decode = execution context storage 5/15/2019 中国科学技术大学 25

26 NVIDIA GeForce GTX 285 core 64 KB of storage for thread contexts (registers) Groups of 32 threads share instruction stream (each group is a Warp) Up to 32 warps are simultaneously interleaved Up to 1024 thread contexts can be stored 5/15/2019 中国科学技术大学

27 NVIDIA GeForce GTX 285 Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex 30 cores on the GTX 285: 30,720 threads 5/15/2019 中国科学技术大学 27

28 GP100 GPU 56 SMs 64 Lanes/SM Or 3584 Stream Processors 5/15/2019 中国科学技术大学 28

29 Summary- 向量处理机 vs. GPU 不同层次相近的术语比较 5/15/2019 中国科学技术大学 29

30 Summary- 向量处理机 vs. GPU 5/15/2019 中国科学技术大学 30

31 Summary-Multimedia SIMD Computers and GPUs 5/15/2019 中国科学技术大学 31

32 5/15/2019 中国科学技术大学 32

33 5/15/2019 中国科学技术大学 33

34 5/15/2019 中国科学技术大学 34

35 GPU Readings Lindholm et al., "NVIDIA Tesla: A Unified Graphics and Computing Architecture," IEEE Micro Fatahalian and Houston, A Closer Look at GPUs, CACM Narasiman et al., Improving GPU Performance via Large Warps and Two-Level Warp Scheduling, MICRO Fung et al., Dynamic Warp Formation and Scheduling for Efficient GPU Control Flow, MICRO Jog et al., Orchestrated Scheduling and Prefetching for GPGPUs, ISCA /15/2019 中国科学技术大学 35

36 Acknowledgements These slides contain material developed and copyright by: John Kubiatowicz (UCB) Krste Asanovic (UCB) John Hennessy (Standford)and David Patterson (UCB) Chenxi Zhang (Tongji) Muhamed Mudawar (KFUPM) UCB material derived from course CS152 CS252 CS61C KFUPM material derived from course COE501 COE502 CMU Introduction to Computer Architecture 5/15/2019 中国科学技术大学 36

投影片 1

投影片 1 2 理 1 2-1 CPU 2-2 CPU 理 2-3 CPU 類 2 什 CPU CPU Central Processing Unit ( 理 ), 理 (Processor), CPU 料 ( 例 ) 邏 ( 例 ),, 若了 CPU, 3 什 CPU CPU 了, 行, 利 CPU 力來行 4 什 CPU 5 2-2-1 CPU CPU 了 (CU, Control Unit) / 邏