add reference and algorith for ch06 (#65)

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`

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<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> 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\*类型指针进行内存读取。
 

mlsys.bib

@@ -497,4 +497,19 @@ pages = {1–9},
 numpages = {9},
 location = {New York, NY, USA},
 series = {ADKDD'14}
-}
+}
+
+@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}
+}