From 937928fcf2dc0e6101e03c9f9c4c8a5bd4bcf7c4 Mon Sep 17 00:00:00 2001 From: Corleone Date: Fri, 11 Mar 2022 16:23:02 +0800 Subject: [PATCH] add reference and algorith for ch06 (#65) --- .../accelerator_architecture.md | 2 +- .../accelerator_programming.md | 34 +- img/ch06/gemm.svg | 2028 ++ img/ch06/gemm_tensor_core.svg | 27334 ++++++++++++++++ mlsys.bib | 17 +- 5 files changed, 29390 insertions(+), 25 deletions(-) create mode 100644 img/ch06/gemm.svg create mode 100644 img/ch06/gemm_tensor_core.svg diff --git a/chapter_accelerator/accelerator_architecture.md b/chapter_accelerator/accelerator_architecture.md index 5bd5a42..e24ed01 100644 --- a/chapter_accelerator/accelerator_architecture.md +++ b/chapter_accelerator/accelerator_architecture.md @@ -5,7 +5,7 @@ ### 硬件加速器的架构 -现代GPU在十分有限的面积上实现了极强的计算能力和极高的储存器以及IO带宽。一块高端的GPU中，晶体管数量已经达到主流CPU的两倍，而且显存已经达到了16GB以上,工作频率也达到了1GHz。GPU的体系架构由两部分组成，分别是流处理阵列和存储器系统，两部分通过一个片上互联网络连接。流处理器阵列和存储器系统都可以单独扩展，规格可以根据产品的市场定位单独裁剪。如GV100的组成如 :numref:`gv100`所示： +现代GPU在十分有限的面积上实现了极强的计算能力和极高的储存器以及IO带宽。一块高端的GPU中，晶体管数量已经达到主流CPU的两倍，而且显存已经达到了16GB以上,工作频率也达到了1GHz。GPU的体系架构由两部分组成，分别是流处理阵列和存储器系统，两部分通过一个片上互联网络连接。流处理器阵列和存储器系统都可以单独扩展，规格可以根据产品的市场定位单独裁剪。如GV100的组成 :cite:`2017NVIDIA`如 :numref:`gv100`所示： ![Volta GV100](../img/ch06/V100.svg) :width:`800px` diff --git a/chapter_accelerator/accelerator_programming.md b/chapter_accelerator/accelerator_programming.md index 0351d08..dbb6c4c 100644 --- a/chapter_accelerator/accelerator_programming.md +++ b/chapter_accelerator/accelerator_programming.md @@ -15,13 +15,11 @@ - **指令层级**：如PTX ISA MMA指令集，提供更细粒度的mma指令，便于用户组成更多种形状的接口，通过CUDA Device端内联编程使能TensorCore。 -矩阵乘法运算作为深度学习网络中占比最大的计算，对其进行优化是十分必要的。因此本节将统一以矩阵乘法$D[M, N] = C[M, N] + A[M, K] * B[K, N]$为实例，对比介绍如何通过不同编程方式使能加速器。 +矩阵乘法运算作为深度学习网络中占比最大的计算，对其进行优化是十分必要的。因此本节将统一以矩阵乘法$D[M, N] = C[M, N] + A[M, K] \times B[K, N]$为实例，对比介绍如何通过不同编程方式使能加速器。 -``` -$A, B$矩阵 $D$矩阵  $C[i][j] \gets 0$  -$C[i][j] \gets C[i][j] + A[i, k] \times B[k, j]$  -$D[i][j] \gets C[i][j]$  -``` +![矩阵乘法GEMM运算](../img/ch06/gemm.svg) +:width:`800px` +:label:`gemm-algorith` ### 硬件加速器的多样化编程方法 :label:`diversified-programming-title` @@ -91,7 +89,7 @@ wmma::fragment c_frag; #### 指令集编程使能加速器 -在NVIDIA PTX ISA (Instruction Set Architecture)中提供了另一个编程接口，如Volta架构中的$mma.sync.m8n8k4$指令，它使用$M=8, N=8, K=4$的形状配置执行乘累加操作。具体地，它由线程组（黑色椭圆表示）或octet执行，如 :numref:`PTX`显示了线程和数据的映射关系。每个线程组由四个连续的线程组成，使用不同颜色的圆圈表示。图中还指出了一个octet里面的线程在线程束内的分布，Float16乘法器A或B的四个连续元素（使用具有相同颜色的块表示），以及Float32累加器C或D的八个分散元素（同样使用相同颜色的块表示）。彩色块上的数字代表对应的线程ID。 +在NVIDIA PTX ISA (Instruction Set Architecture)中提供了另一个编程接口，如Volta架构中的$mma.sync.m8n8k4$指令，它使用$M=8, N=8, K=4$的形状配置执行乘累加操作。具体地，它由线程组（黑色椭圆表示）或octet执行 :cite:`2018Modeling`，如 :numref:`PTX`显示了线程和数据的映射关系。每个线程组由四个连续的线程组成，使用不同颜色的圆圈表示。图中还指出了一个octet里面的线程在线程束内的分布，Float16乘法器A或B的四个连续元素（使用具有相同颜色的块表示），以及Float32累加器C或D的八个分散元素（同样使用相同颜色的块表示）。彩色块上的数字代表对应的线程ID。 ![mma指令之线程与矩阵元素映射关系](../img/ch06/ptx.svg) :width:`800px` @@ -123,21 +121,11 @@ res = te.lang.cce.matmul(tensor_a, tensor_b, False, False, False, dst_dtype=dst_ 本节 :numref:`accelerator-program-title`前几个小节主要介绍了硬件加速器的不同层级的多样化编程方法。调用计算库的方式留给程序员的优化空间较少，合理利用硬件加速器不同层级的编程，可以实现更好的性能优化。 为了更好的让读者理解硬件加速器的使用，本节会继续 :numref:`accelerator-programable-title`节中的GEMM运算，仍以WMMA API使能Tensor Core加速单元为例，介绍如何通过矩阵分块、资源映射等方式更高效的利用硬件加速器。 -[\[alg:TensorCore\]]{#alg:TensorCore label="alg:TensorCore"} +![TensorCore矩阵乘法GEMM运算](../img/ch06/gemm_tensor_core.svg) +:width:`800px` +:label:`gemm-tensor-core-algorith` -``` -$A, B$矩阵 $D$矩阵   Mapping to Block.Idx  Mapping to Block.Idy  Mapping -to Block.Idz -$A_{Shared}[i_o][k_o] \gets A[i_o][k_o]$ $B_{Shared}[k_o][j_o] \gets B[k_o][j_o]$ $Syncthreads()$  - Mapping to Warp.Idx  Mapping to Warp.Idy -$A_{Register}[i_i][k_i] \gets A_{Shared}[i_i][k_i]$ $B_{Register}[k_i][j_i] \gets B_{Shared}[k_i][j_i]$ $pragma\ unroll$  -$wmma.load\_matrix\_sync(A_{Fragment}, A_{Register})$ $wmma.load\_matrix\_sync(B_{Fragment}, B_{Register})$  -$wmma.fill\_fragment(C_{Fragment}, 0)$  -$wmma.mma\_sync(D_{Fragment}, C_{Fragment}, A_{Fragment}, B_{Fragment})$  -$Syncthreads()$  $wmma.store\_matrix\_sync(D, D_{Fragment})$  -``` - -若要得到高性能CUDA程序，提高并行性、增大吞吐量、优化指令执行是至关重要的三个优化目标。针对该实例，具体地实现和优化方案列出如下，对应到具体实例伪代码如算法2所示： +若要得到高性能CUDA程序，提高并行性、增大吞吐量、优化指令执行是至关重要的三个优化目标。针对该实例，具体地实现和优化方案列出如下，对应到具体实例伪代码如 :numref:`gemm-tensor-core-algorith`所示： 1. **优化内存结构------增大吞吐量**：将原始大规模矩阵根据不同阈值切分成不同层级的子矩阵块，使得子矩阵块能被如共享内存、寄存器等高性能体系结构存储下来，以此提高吞吐量。设置切分参数为$BlockTile[Ms, Ns, Ks]$和$WarpTile[Mw, Nw, Kw]$，对应的将BlockTile下的矩阵由全局内存搬移至共享内存，以提高全局内存合并访问和数据局部性，如 :numref:`GEMM-BlockTile`所示；再将内层WarpTile下的矩阵由共享内存搬移至寄存器中，如 :numref:`GEMM-WarpTile`所示，以备Tensor Core加速器数据存取。 @@ -149,7 +137,7 @@ $Syncthreads()$  $wmma.store\_matrix\_sync(D, D_{Fragment})$  :width:`800px` :label:`GEMM-WarpTile` -2. **并行资源映射------提高并行性**：将多层级的并行资源（Block、Warp、Thread）与对应需要计算/搬移的数据建立映射关系，提高程序并行性。将可并行的计算/数据搬移操作映射到并行资源上，对于GEMM实例，M/N轴即为可并行轴，将数据搬移操作中的循环指令映射分配到Block层级（即算法3中的2-4行$For$循环），将内层循环指令映射分配到Warp层级（即算法3中的8-10行$For$循环）。（前文介绍，线程束Warp作为调度的基本单位，且是WMMA API操纵的基本层级，因此对Warp层级进行数据映射比Thread层级映射更为合适） +2. **并行资源映射------提高并行性**：将多层级的并行资源（Block、Warp、Thread）与对应需要计算/搬移的数据建立映射关系，提高程序并行性。将可并行的计算/数据搬移操作映射到并行资源上，对于GEMM实例，M/N轴即为可并行轴，将数据搬移操作中的循环指令映射分配到Block层级（即 :numref:`gemm-tensor-core-algorith`中的2-4行$For$循环），将内层循环指令映射分配到Warp层级（即 :numref:`gemm-tensor-core-algorith`中的8-9行$For$循环）。（前文介绍，线程束Warp作为调度的基本单位，且是WMMA API操纵的基本层级，因此对Warp层级进行数据映射比Thread层级映射更为合适） 3. **Warp统一的Tensor Core数据交互------增大吞吐量**：根据 :numref:`diversified-programming-title`节中介绍的编程方法，除调用算子库外，均需要使用或将指令封装成WMMA接口形式统一进行Warp层级的数据存取和计算。如 :numref:`GEMM-TensorCore`所示，Tensor Core加速器需要从局部内存/寄存器中读取数据，存于虚拟Fragment数据结构中，对应使用$wmma.load\_matrix\_sync()$接口，将累加Fragment $C$ 通过$wmma.fill\_fragment()$接口进行初始化后，使用$wmma.mma\_sync()$使能加速器进行乘累加运算，后将结果Fragment $D$通过调用$wmma.store\_matrix\_sync()$接口拷贝至目标内存地址。 @@ -161,5 +149,5 @@ $Syncthreads()$  $wmma.store\_matrix\_sync(D, D_{Fragment})$  5. **资源负载均衡------增大吞吐量**：调整平衡每个线程处理的数据量、共享内存使用量、寄存器使用量，以获得更高的SM占用率。一般在实际程序中BlockTile和WarpTile的选取至关重要。 -6. **优化指令执行**：使用\#unroll功能进行循环展开以避免分支冲突，如算法3中13行；使用向量化加载指令减少PTX指令执行次数以提高带宽等，对于GPU Volta架构，最大向量化加载指令为ldg128，即128比特带宽，对于算法3中5-6行数据由全局内存加载至共享内存时，即可采用Float4\*类型指针进行内存读取。 +6. **优化指令执行**：使用\#unroll功能进行循环展开来提升指令级并行，如 :numref:`gemm-tensor-core-algorith`中13行；使用向量化加载指令以提高带宽等，对于GPU Volta架构，最大向量化加载指令为ldg128，即128比特带宽，对于 :numref:`gemm-tensor-core-algorith`中5-6行数据由全局内存加载至共享内存时，即可采用Float4\*类型指针进行内存读取。 diff --git a/img/ch06/gemm.svg b/img/ch06/gemm.svg new file mode 100644 index 0000000..822818e --- /dev/null +++ b/img/ch06/gemm.svg @@ -0,0 +1,2028 @@ + + + + + + + + + + + + + + diff --git a/img/ch06/gemm_tensor_core.svg b/img/ch06/gemm_tensor_core.svg new file mode 100644 index 0000000..0a4183f --- /dev/null +++ b/img/ch06/gemm_tensor_core.svg @@ -0,0 +1,27334 @@ + + + + + + + + + + + + + + diff --git a/mlsys.bib b/mlsys.bib index a330def..660795d 100644 --- a/mlsys.bib +++ b/mlsys.bib @@ -497,4 +497,19 @@ pages = {1–9}, numpages = {9}, location = {New York, NY, USA}, series = {ADKDD'14} -} \ No newline at end of file +} + +@misc{2017NVIDIA, + author={NVIDIA}, + title={NVIDIA Tesla V100 GPU Architecture: The World's Most Advanced Datacenter GPU}, + year={2017}, + howpublished = "Website", + note = {\url{http://www.nvidia.com/object/volta-architecture-whitepaper.html}} +} + +@article{2018Modeling, + title={Modeling Deep Learning Accelerator Enabled GPUs}, + author={Raihan, M. A. and Goli, N. and Aamodt, T.}, + journal={arXiv e-prints arXiv:1811.08309}, + year={2018} +}