sync overleaf for acceleerator (#420)

Co-authored-by: Corleone <liuchao195@huawei.com> Co-authored-by: Jiarong Han <jiaronghan@outlook.com>
2026-04-08 05:09:17 +08:00 · 2023-02-17 10:16:44 +08:00
parent 0c0bff1b83
commit 039787dc8d
4 changed files with 71 additions and 56 deletions
--- a/chapter_accelerator/accelerator_architecture.md
+++ b/chapter_accelerator/accelerator_architecture.md
@@ -1,10 +1,10 @@
 ## 加速器基本组成原理

-上节主要介绍了加速器的意义以及设计思路，了解到加速器与通用处理器在设计上的区别，因此加速器的硬件结构与CPU的硬件结构有着根本的不同，通常都是由多种片上缓存以及多种运算单元组成。本章节主要通过GPU的Volta架构作为样例进行介绍。
+上节主要介绍了加速器的意义以及设计思路，讲述了加速器与通用处理器在设计上的区别，可以看到加速器的硬件结构与CPU的硬件结构有着根本的不同，通常都是由多种片上缓存以及多种运算单元组成。本章节主要以GPU的Volta架构作为样例进行介绍。

 ### 硬件加速器的架构

-现代GPU在十分有限的面积上实现了极强的计算能力和极高的储存器以及IO带宽。一块高端的GPU中，晶体管数量已经达到主流CPU的两倍，而且显存已经达到了16GB以上,工作频率也达到了1GHz。GPU的体系架构由两部分组成，分别是流处理阵列和存储器系统，两部分通过一个片上互联网络连接。流处理器阵列和存储器系统都可以单独扩展，规格可以根据产品的市场定位单独裁剪。如GV100的组成 :cite:`2017NVIDIA`如 :numref:`gv100`所示：
+现代GPU在十分有限的面积上实现了极强的计算能力和极高的储存器以及IO带宽。在一块高端的GPU中，晶体管数量已经达到主流CPU的两倍，而且显存已经达到了16GB以上，工作频率也达到了1GHz。GPU的体系架构由两部分组成，分别是流处理阵列和存储器系统，两部分通过一个片上互联网络连接。流处理器阵列和存储器系统都可以单独扩展，规格可以根据产品的市场定位单独裁剪。如GV100的组成 :cite:`2017NVIDIA`如 :numref:`gv100`所示：

 ![Volta GV100](../img/ch06/V100.svg)
 :width:`800px`
@@ -17,32 +17,31 @@
    -   64个32位浮点运算单元
    -   64个32位整数运算单元
    -   32个64位浮点运算单元
-    -   8个张量核
+    -   8个张量计算核心
    -   4个纹理单元
-   8个512-bit内存控制器
+-   8个512位内存控制器

 一个完整的GV100 GPU含有84个SM，5376个32位浮点运算单元，5376个32位整型运算单元，2688个64位浮点运算单元，672个张量运算单元和336个纹理单元。一对内存控制器控制一个HBM2 DRAM堆栈。 :numref:`gv100`中展示的为带有84个SM的GV100 GPU(不同的厂商可以使用不同的配置)，Tesla V100则含有80个SM。

 ### 硬件加速器的存储单元

-与传统的CPU模型相似，从一个计算机系统主内存DRAM中获取数据的速度相对于处理器的运算速度较慢。对于加速器而言，如果没有缓存进行快速存取，DRAM的带宽非常不足。如果无法快速的在DRAM上获取程序和数据，加速器将因空置而降低利用率。为了缓解DRAM的带宽问题，GPU提供了不同层次的若干区域供程序员存放数据，每块区域的内存都有自己的最大带宽以及延迟。
+与传统的CPU模型相似，从一个计算机系统主内存DRAM中获取数据的速度相对于处理器的运算速度较慢。对于加速器而言，如果没有缓存进行快速存取，DRAM的带宽非常不足。如果无法快速地在DRAM上获取程序和数据，加速器将因空置而降低利用率。为了缓解DRAM的带宽问题，GPU提供了不同层次的若干区域供程序员存放数据，每块区域的内存都有自己的最大带宽以及延迟。开发者需根据不同存储器之间的存储速度的数量级的变化规律，选用适当类型的内存以及最大化地利用它们，从而发挥硬件的最大算力，减少计算时间。

-   **寄存器文件（Register File）**：片上最快的存储器，但与CPU不同，GPU的每个SM（流多处理器）有上万个寄存器。但当每个线程使用过多的寄存器时，SM中能够调度的线程块数量就会受到限制，可执行的线程总数量会因此受到限制，可执行的线程数量过少会造成硬件无法充分的利用，性能急剧下降。所以要根据算法的需求合理使用寄存器。
-
-   **共享内存（Shared Memory）**：共享内存实际上是用户可控的一级缓存，每个SM（流多处理器）中有128KB的一级缓存, 开发者可根据应用程序需要配置最大96KB的一级缓存作为共享内存。共享内存的访存延迟极低，只有几十个时钟周期。共享内存具有高达1.5TB/s的带宽，远远高于全局内存的峰值带宽900GB/s。所以说，共享内存的使用对于一个高性能计算工程师来说是一个必须要掌握的一个概念。
+-   **寄存器文件（Register File）**：片上最快的存储器，但与CPU不同，GPU的每个SM（流多处理器）有上万个寄存器。尽管如此当每个线程使用过多的寄存器时，SM中能够调度的线程块数量就会受到限制，可执行的线程总数量会因此受到限制，可执行的线程数量过少会造成硬件无法充分的利用，性能急剧下降。所以要根据算法的需求合理使用寄存器。
+ 
+-   **共享内存（Shared Memory）**：共享内存实际上是用户可控的一级缓存，每个SM（流多处理器）中有128KB的一级缓存, 开发者可根据应用程序需要配置最大96KB的一级缓存作为共享内存。共享内存的访存延迟极低，只有几十个时钟周期。共享内存具有高达1.5TB/s的带宽，远远高于全局内存的峰值带宽900GB/s。共享内存的使用对于高性能计算工程师来说是一个必须要掌握的概念。

 -   **全局内存（Global Memory）**：全局内存之所以称为全局，是因为GPU与CPU都可以对它进行读写操作。全局内存对于GPU中的每个线程都是可见的，都可以直接对全局内存进行读写操作。CPU等其他设备可以通过PCI-E总线对其进行读写操作。全局内存也是GPU中容量最大的一块内存，可达16GB之多。同时也是延迟最大的内存，通常有高达上百个时钟周期的访存延迟。

-   **常量内存（Constant Memory）**：常量内存其实只是全局内存的一种虚拟地址形式，并没有真正的物理硬件内存块。常量内存有两个特性，一个高速缓存，另一个更重要的特性是它支持将某个单个值广播到线程束中的每个线程中。
+-   **常量内存（Constant Memory）**：常量内存其实只是全局内存的一种虚拟地址形式，并没有真正的物理硬件内存块。常量内存有两个特性，一个是高速缓存，另一个更重要的特性是它支持将某个单个值广播到线程束中的每个线程中。

-   **纹理内存（Texture Memory）**：纹理内存是全局内存的一个特殊形态。当全局内存被绑定为纹理内存时，执行读写操作将通过专用的纹理缓存来加速。在早期的GPU上没有缓存，因此每个SM（流多处理器）上的纹理内存为设备提供了唯一真正缓存数据的方法。纹理内存的另外一个特性，也是最有用的特性就是当访问存储单元时，允许GPU实现硬件相关的操作。比如说使用纹理内存，可以通过归一化的地址对数组进行访问，获取的数据可以通过硬件进行自动插值，从而达到快速处理数据的目的。此外对于二维数组和三维数组，支持硬件级的双线性插值与三线性插值。纹理内存另一个实用的特性是可以根据数组的索引自动处理边界条件，不需要对特殊边缘进行处理即可完成数组内元素操作，从而防止线程中分支的产生。
+-   **纹理内存（Texture Memory）**：纹理内存是全局内存的一个特殊形态。当全局内存被绑定为纹理内存时，执行读写操作将通过专用的纹理缓存来加速。在早期的GPU上没有缓存，因此每个SM上的纹理内存为设备提供了唯一真正缓存数据的方法。然而随着硬件的升级，一级缓存和二级缓存的出现使得纹理缓存的这项优势已经荡然无存。纹理内存的另外一个特性，也是最有用的特性就是当访问存储单元时，允许GPU实现硬件相关的操作。比如说使用纹理内存，可以通过归一化的地址对数组进行访问，获取的数据可以通过硬件进行自动插值，从而达到快速处理数据的目的。此外对于二维数组和三维数组，支持硬件级的双线性插值与三线性插值。纹理内存另一个实用的特性是可以根据数组的索引自动处理边界条件，不需要对特殊边缘进行处理即可完成数组内元素操作，从而防止线程中分支的产生。

-由于寄存器的高速读取特性，因此每次计算都离不开寄存器的参与。接着是一级缓存和共享内存，然后是常量内存、纹理内存、全局内存，最后则是主机端内存。根据不同存储器之间的存储速度的数量级的变化规律，选用适当类型的内存以及最大化地利用它们，从而发挥硬件的最大算力，减少计算时间。

 ### 硬件加速器的计算单元
 :label:`compute-unit-title`

-为了支持不同的神经网络模型，加速器会提供以下几种计算单元，不同的网络层可以根据需要选择使用对应的计算单元。如 :numref:`compute-unit`所示
+为了支持不同的神经网络模型，加速器会提供以下几种计算单元，不同的网络层可以根据需要选择使用合适的计算单元，如 :numref:`compute-unit`所示

 -   **标量计算单元**：与标准的精简指令运算集（Reduced Instruction Set Computer，RISC）相似，一次计算一个标量元素。

@@ -56,30 +55,31 @@
 :width:`800px`
 :label:`compute-unit`

-GPU计算单元主要由标量计算单元组成，而在Volta及以后的架构中还加入了三维向量计算单元。如 :numref:`SM`所示,对于每个SM，其中64个32位浮点运算单元、64个32位整数运算单元、32个64位浮点运算单元均为标量计算单元。而8个张量核则是专为神经网络应用设计的三维向量计算单元。
+GPU计算单元主要由标量计算单元和三维向量计算单元组成。如 :numref:`SM`所示,对于每个SM，其中64个32位浮点运算单元、64个32位整数运算单元、32个64位浮点运算单元均为标量计算单元。而8个张量计算核心则是专为神经网络应用设计的三维向量计算单元。

 ![Volta GV100 流多处理器（SM）](../img/ch06/SM.svg)
 :width:`800px`
 :label:`SM`

-张量核（Tensor Core）每个时钟周期完成一次$4\times4$的矩阵乘累加计算,如 :numref:`tensorcore`：
+张量计算核心每个时钟周期完成一次$4\times4$的矩阵乘累加计算,如 :numref:`tensorcore`所示：

 ```cpp
    D = A * B + C
 ```

-![Tensor Core $4\times4$矩阵乘累加计算](../img/ch06/tensor_core.svg)
+![张量计算核心$4\times4$矩阵乘累加计算](../img/ch06/tensor_core.svg)
 :width:`800px`
 :label:`tensorcore`

-其中A,B,C和D都是$4\times4$的矩阵，矩阵乘累加的输入矩阵A和B是FP16的矩阵，累加矩阵C和D可以是FP16也可以是FP32。 V100的张量核是可编程的矩阵乘法和累加计算单元，可以提供多达125 Tensor TFLOPS(Tera Floating-point Operations Per Second)的训练和推理应用。相比于普通的FP32计算单元可以提速10倍以上。
+其中A,B,C和D都是$4\times4$的矩阵，矩阵乘累加的输入矩阵A和B是FP16的矩阵，累加矩阵C和D可以是FP16也可以是FP32。 V100的张量计算核心是可编程的矩阵乘法和累加计算单元，可以提供多达125 Tensor TFLOPS(Tera Floating-point Operations Per Second)的训练和推理应用。相比于普通的FP32计算单元可以提速10倍以上。

 ### DSA芯片架构

-为了满足飞速发展的深度神经网络对芯片算力的需求，业界也纷纷推出了特定领域架构DSA芯片设计。以华为公司昇腾系列AI处理器为例，本质上是一个片上系统（System on Chip，SoC），主要应用在图像、视频、语音、文字处理相关的场景。主要的架构组成部件包括特制的计算单元、大容量的存储单元和相应的控制单元。该芯片由以下几个部分构成：芯片系统控制CPU（Control CPU），AI计算引擎（包括AI Core和AI CPU），多层级的片上系统缓存（Cache）或缓冲区（Buffer），数字视觉预处理模块（Digital Vision Pre-Processing，DVPP）等。
+为了满足飞速发展的深度神经网络对芯片算力的需求，业界也纷纷推出了特定领域架构DSA芯片设计。以华为公司昇腾系列AI处理器为例，本质上是一个片上系统（System on Chip，SoC），主要应用在图像、视频、语音、文字处理相关的场景。主要的架构组成部件包括特制的计算单元、大容量的存储单元和相应的控制单元。该芯片由以下几个部分构成：芯片系统控制CPU（Control CPU）、AI计算引擎（包括AI Core和AI CPU）、多层级的片上系统缓存（Cache）或缓冲区（Buffer）、数字视觉预处理模块（Digital Vision Pre-Processing，DVPP）等。

 ![达芬奇架构设计](../img/ch06/davinci_architecture.svg)
 :width:`800px`
 :label:`davinci_architecture`

-昇腾AI芯片的计算核心主要由AI Core构成，负责执行标量、向量和张量相关的计算密集型算子。AI Core采用了达芬奇架构 :cite:`2021Ascend`，基本结构如 :numref:`davinci_architecture`所示，从控制上可以看成是一个相对简化的现代微处理器基本架构。它包括了三种基础计算单元：矩阵计算单元（Cube Unit）、向量计算单元（Vector Unit）和标量计算单元（Scalar Unit）。这三种计算单元分别对应了张量、向量和标量三种常见的计算模式，在实际的计算过程中各司其职，形成了三条独立的执行流水线，在系统软件的统一调度下互相配合达到优化计算效率的目的。 同GPU类似，在矩阵乘加速设计上，在AICore中也提供了矩阵计算单元作为昇腾AI芯片的核心计算模块，意图高效解决矩阵计算的瓶颈问题。矩阵计算单元提供强大的并行乘加计算能力，可以用一条指令完成两个$16\times16$矩阵的相乘运算，等同于在极短时间内进行了$16\times16\times16=4096$个乘加运算，并且可以实现FP16的运算精度。
+昇腾AI芯片的计算核心主要由AI Core构成，负责执行标量、向量和张量相关的计算密集型算子。AI Core采用了达芬奇架构 :cite:`2021Ascend`，基本结构如 :numref:`davinci_architecture`所示，从控制上可以看成是一个相对简化的现代微处理器基本架构。它包括了三种基础计算单元：矩阵计算单元（Cube Unit）、向量计算单元（Vector Unit）和标量计算单元（Scalar Unit）。这三种计算单元分别对应了张量、向量和标量三种常见的计算模式，在实际的计算过程中各司其职，形成了三条独立的执行流水线，在系统软件的统一调度下互相配合达到优化计算效率的目的。
+同GPU类似，在矩阵乘加速设计上，在AICore中也提供了矩阵计算单元作为昇腾AI芯片的核心计算模块，意图高效解决矩阵计算的瓶颈问题。矩阵计算单元提供强大的并行乘加计算能力，可以用一条指令完成两个$16\times16$矩阵的相乘运算，等同于在极短时间内进行了$16\times16\times16=4096$个乘加运算，并且可以实现FP16的运算精度。
--- a/chapter_accelerator/accelerator_introduction.md
+++ b/chapter_accelerator/accelerator_introduction.md
@@ -2,17 +2,16 @@

 ### 硬件加速器设计的意义

-未来人工智能发展的三大核心要素是数据、算法和算力。目前，人工智能系统算力大都构建在CPU+GPU之上，主体多是GPU。随着神经网络层数的增多，模型体量的增大，算法复杂度的上升，CPU和GPU很难再满足新型网络对于算力的需求。例如，2015年谷歌的AlphaGo与[樊麾](https://baike.baidu.com/item/樊麾)对弈时，用了1202个CPU和176个GPU，每盘棋需要消耗上千美元的电费，而与之对应的是樊麾的功耗仅为20瓦。
+未来人工智能发展的三大核心要素是数据、算法和算力。目前，人工智能系统算力大都构建在CPU和GPU之上且主体多是GPU。随着神经网络的层增多，模型体量增大，算法趋于复杂，CPU和GPU很难再满足新型网络对于算力的需求。例如，2015年谷歌的AlphaGo用了1202个CPU和176个GPU打败了人类职业选手，每盘棋需要消耗上千美元的电费，而与之对应的是人类选手的功耗仅为20瓦。
+
+虽然GPU在面向向量、矩阵以及张量的计算上，引入许多新颖的优化设计，但由于GPU需要支持的计算类型复杂，芯片规模大、能耗高，人们开始将更多的精力转移到深度学习硬件加速器的设计上来。和传统CPU和GPU芯片相比，深度学习硬件加速器有更高的性能和更低的能耗。未来随着人们真正进入智能时代，智能应用的普及会越来越广泛，到那时每台服务器、每台智能手机和每个智能摄像头，都需要使用深度学习加速器。

-虽然GPU在面向向量、矩阵以及张量的计算上，引入许多新颖的优化设计，但由于GPU需要支持的计算类型复杂，芯片规模大、能耗高，人们开始将更多的精力转移到深度学习硬件加速器的设计上来。和传统CPU和GPU芯片相比，新型深度学习加速器会有更高的性能，以及更低的能耗。未来随着人们真正进入智能时代，智能应用的普及会越来越广泛，到那时每台服务器、每台智能手机、每个智能摄像头，都需要使用加速器。

 ### 硬件加速器设计的思路
 :label:`accelerator-design-title`

-近些年来，计算机体系结构的研究热点之一就是深度学习硬件加速器的设计。在体系结构的研究中，能效和通用性是两个重要的衡量指标。能效关注单位能耗下基本计算的次数，通用性主要指芯片能够覆盖的任务种类。
+近些年来，计算机体系结构的研究热点之一是深度学习硬件加速器的设计。在体系结构的研究中，能效和通用性是两个重要的衡量指标。其中能效关注单位能耗下基本计算的次数，通用性主要指芯片能够覆盖的任务种类。以两类特殊的芯片为例：一种是较为通用的通用处理器(如CPU)，该类芯片理论上可以完成各种计算任务，但是其能效较低大约只有0.1TOPS/W。另一种是专用集成电路(Application Specific Integrated Circuit, ASIC)，其能效更高，但是支持的任务相对而言就比较单一。对于通用的处理器而言，为了提升能效，在芯片设计上引入了许多加速技术，例如：超标量技术、单指令多数据（Single Instruction Multiple Data，SIMD）技术以及单指令多线程（Single Instruction Multiple Threads，SIMT）技术等。

-以两类特殊的芯片为例:一种是我们较为熟悉的通用处理器(如CPU)，该类芯片理论上可以完成各种计算任务，但是其能效较低大约只有0.1TOPS/W；另一种是专用集成电路(Application Specific Integrated Circuit, ASIC)，其能效更高，但是支持的任务相对而言就比较单一。对于通用的处理器而言，为了提升能效，在芯片设计上有许多加速技术的引入，例如：超标量技术、单指令多数据（Single Instruction Multiple Data，SIMD）技术以及单指令多线程（Single Instruction Multiple Threads，SIMT）技术等。
+对于不同的加速器设计方向，业界也有不同的硬件实现。针对架构的通用性，NVIDIA持续在GPU芯片上发力，先后推出了Volta、 Turing、 Ampere等架构，并推出用于加速矩阵计算的张量计算核心（Tensor Core），以满足深度学习海量算力的需求。

-对于不同的加速器设计方向，业界也有不同的硬件实现。针对架构的通用性，NVIDIA持续在其GPU芯片上发力，先后推出了Volta, Turing, Ampere架构，并推出用于加速矩阵计算的张量核（Tensor Core），以满足深度学习海量算力的需求。
-
-对于偏定制化的硬件架构，面向深度学习计算任务，业界提出了特定领域架构(Domain Specific Architecture)。 Google公司推出了TPU芯片，专门用于加速深度学习计算任务，其使用脉动阵列(Systolic Array)来优化矩阵乘法和卷积运算，可以充分地利用数据局部性，降低对内存的访问次数。华为也推出了自研的昇腾AI处理器，旨在为用户提供更高能效的算力和易用的开发、部署体验，其中的CUBE运算单元，就用于加速矩阵乘法的计算。
+对于偏定制化的硬件架构，面向深度学习计算任务，业界提出了特定领域架构(Domain Specific Architecture， DSA)。Google公司推出了TPU芯片，专门用于加速深度学习计算任务，其使用脉动阵列(Systolic Array)来优化矩阵乘法和卷积运算，可以充分地利用数据局部性，降低对内存的访问次数。华为也推出了自研昇腾AI处理器，旨在为用户提供更高能效的算力和易用的开发、部署体验，其中的CUBE运算单元，就用于加速矩阵乘法的计算。
--- a/chapter_accelerator/accelerator_programming.md
+++ b/chapter_accelerator/accelerator_programming.md
@@ -1,7 +1,7 @@
 ## 加速器基本编程原理
 :label:`accelerator-program-title`

-本章前两节主要介绍了硬件加速器设计的意义、思路以及基本组成原理。软硬件协同优化作为构建高效AI系统的一个重要指导思想，需要软件算法/软件栈和硬件架构在神经网络应用中互相影响、紧密耦合。为了最大限度地发挥加速器的优势，要求能够基于硬件系统架构提供易用、高效的编程方法。因此，在本节中将着重介绍加速器的可编程性，包括编程接口直接调用方式及算子编译器优化方式。
+本章前两节主要介绍了这些硬件加速器设计的意义、思路以及基本组成原理。软硬件协同优化作为构建高效AI系统的一个重要指导思想，需要软件算法/软件栈和硬件架构在神经网络应用中互相影响、紧密耦合。为了最大限度地发挥加速器的优势，要求能够基于硬件系统架构设计出一套较为匹配的指令或编程方法。因此，本节将着重介绍加速器的可编程性，以及如何通过编程使能加速器，提升神经网络算子的计算效率。

 ### 硬件加速器的可编程性
 :label:`accelerator-programable-title`
@@ -10,25 +10,25 @@

 #### 编程接口使能加速器

-硬件加速器出于计算效率和易用性等方面考虑，将编程接口使能方式分为不同等级，一般包括：算子库层级，编程原语层级，以及指令层级。为了更具象的解释上述层级的区别，我们以Volta架构的Tensor Core加速器为例，由高层至底层对比介绍这三种不同编程方式：
+硬件加速器出于计算效率和易用性等方面考虑，将编程使能方式分为不同等级，一般包括：算子库层级，编程原语层级，以及指令层级。为了更具象的解释上述层级的区别，仍以Volta架构的张量计算核心为例，由高层至底层对比介绍这三种不同编程方式：

-   **算子库层级**：如cuBLAS基本矩阵与向量运算库，cuDNN深度学习加速库，均通过Host端调用算子库提供的核函数使能TensorCore；
+-   **算子库层级**：如cuBLAS基本矩阵与向量运算库，cuDNN深度学习加速库，均通过Host端调用算子库提供的核函数使能张量计算核心；

-   **编程原语层级**：如基于CUDA的WMMA API编程接口。同算子库相比，需要用户显式调用计算各流程，如矩阵存取至TensorCore、TensorCore执行矩阵乘累加运算、TensorCore累加矩阵数据初始化操作等；
+-   **编程原语层级**：如基于CUDA的WMMA API编程接口。同算子库相比，需要用户显式调用计算各流程，如矩阵存取至寄存器、张量计算核心执行矩阵乘累加运算、张量计算核心累加矩阵数据初始化操作等；

-   **指令层级**：如PTX ISA MMA指令集，提供更细粒度的mma指令，便于用户组成更多种形状的接口，通过CUDA Device端内联编程使能TensorCore。
+-   **指令层级**：如PTX ISA MMA指令集，提供更细粒度的mma指令，便于用户组成更多种形状的接口，通过CUDA Device端内联编程使能张量计算核心。

 #### 算子编译器使能加速器

- DSA架构的多维度AI加速器通常提供了更多的指令选择（3D-Matrix/2D-Vector/1D-Scalar），以及更加复杂的数据流处理，通过提供接口调用的方式对程序开发人员带来较大的挑战。此外，由于调度、切分的复杂度增加，直接提供算子库的方式由于缺少根据目标shape调优的能力，往往无法在所有shape下均得到最优的性能。因此，对于DSA加速器，业界通常采用算子编译器的解决方案。
+DSA架构的多维度AI加速器通常提供了更多的指令选择（三维向量计算指令、二维向量计算指令、一维向量计算指令），以及更加复杂的数据流处理，通过提供接口调用的方式对程序开发人员带来较大的挑战。此外，由于调度、切分的复杂度增加，直接提供算子库的方式由于缺少根据目标形状（Shape）调优的能力，往往无法在所有形状下均得到最优的性能。因此，对于DSA加速器，业界通常采用算子编译器的解决方案。

- 随着深度学习模型的迭代更新及各类AI芯片的层出不穷，基于人工优化算子的方式给算子开发团队带来沉重的负担。因此，开发一种能够将High-level的算子表示编译成目标硬件可执行代码的算子编译器，逐渐成为学术界及工业界的共识。算子编译器前端通常提供了特定领域描述语言（DSL），用于定义算子的计算范式；类似于传统编译器，算子编译器也会将算子计算表示转换为中间表示，如HalideIR :cite:`ragan2013halide`、TVM :cite:`chen2018tvm`的TIR、Schedule Tree :cite:`verdoolaege2010isl`等，基于模板（手动）、搜索算法或优化求解算法（自动）等方式完成循环变换、循环切分等调度相关优化，以及硬件指令映射、内存分配、指令流水等后端pass优化，最后通过codegen模块将IR转换为DSA加速器可执行的kernel。
+ 随着深度学习模型的迭代更新及各类AI芯片的层出不穷，基于人工优化算子的方式给算子开发团队带来沉重的负担。因此，开发一种能够将High-level的算子表示编译成目标硬件可执行代码的算子编译器，逐渐成为学术界及工业界的共识。算子编译器前端通常提供了特定领域描述语言（DSL），用于定义算子的计算范式；类似于传统编译器，算子编译器也会将算子计算表示转换为中间表示，如HalideIR :cite:`ragan2013halide`、TVM :cite:`chen2018tvm`的TIR、Schedule Tree :cite:`verdoolaege2010isl`等，基于模板（手动）、搜索算法或优化求解算法（自动）等方式完成循环变换、循环切分等调度相关优化，以及硬件指令映射、内存分配、指令流水等后端pass优化，最后通过代码生成模块将IR转换为DSA加速器可执行的设备端核函数。

- 当前业界的算子编译器/编译框架主要有TVM/Ansor :cite:`zheng2020ansor`、MLIR :cite:`lattner2020mlir`、以及华为Ascend芯片上的TBE/AKG :cite:`zhao2021akg`等。
+ 当前业界的算子编译器/编译框架主要有TVM/Ansor :cite:`zheng2020ansor`、MLIR :cite:`lattner2020mlir`、以及华为昇腾芯片上的TBE/AKG :cite:`zhao2021akg`等。

 -   **TVM/Ansor**
 
- TVM是陈天奇博士等人开发的开源深度学习编译框架，提供了端到端的编译优化（图优化/算子优化）能力，在工业界应用较广。在架构上，主要包括relay和tir两层。通过relay导入推理模型，进行算子融合等图层优化，通过tir生成融合算子。在算子编译方面，TVM采用了计算和调度分离的技术，为不同的算子提供了不同的模板，同时支持自定义模板，优化特定算子类型调度。为了更进一步优化算子性能，TVM支持对算子进行自动tuning，来生成较优的切分参数。此外，为了简化用户开发模板的工作，TVM在0.8版本后提供了自动调度能力Ansor，通过搜索的方式，为目标算子生成调度及切分参数。
+TVM是陈天奇博士等人开发的开源深度学习编译框架，提供了端到端的编译优化（图优化/算子优化）能力，在工业界应用较广。在架构上，主要包括Relay和TIR两层。通过Relay导入推理模型，进行算子融合等图层优化，通过TIR生成融合算子。在算子编译方面，TVM采用了计算和调度分离的技术，为不同的算子提供了不同的模板，同时支持自定义模板，优化特定算子类型调度。为了更进一步优化算子性能，TVM支持对算子进行自动调优，来生成较优的切分参数。此外，为了简化用户开发模板的工作，TVM在0.8版本后提供了自动调度能力Ansor，通过搜索的方式，为目标算子生成调度及切分参数。

 ![TVM](../img/ch06/TVM.svg)
 :width:`800px`
@@ -37,7 +37,7 @@

 -   **MLIR**

- 前面的章节介绍过，Google开发的MLIR并不是一个单一的算子编译器，而是一套编译器基础设施，提供了工具链的组合与复用能力。基于MLIR，DSA加速器厂商可以快速的搭建其定制化算子编译器。如Google论文 :cite:`vasilache2022composable`中所述，当前的算子编译器大多提供了一整套自顶向下的编译优化pass，包括调度优化、切分优化、窥孔优化、后端优化、指令生成等，彼此之间大多无法复用，导致新的场景中通常又得从头开发。而在MLIR中，将功能相近的IR优化pass封装为方言（Dialect），并且提供了多个代码生成相关的基础方言，如vector、memref、tensor、scf、affine、linalg等。硬件厂商可以基于这些Dialect，快速构建一整套lower优化及codegen流程。如 :numref:`MLIR_Lowing`所示，利用scf、affine、linalg等方言，对结构化的计算IR完成循环并行优化、切分、向量化等，最后基于LLVM完成指令映射。
+前面的章节介绍过，Google开发的MLIR并不是一个单一的算子编译器，而是一套编译器基础设施，提供了工具链的组合与复用能力。基于MLIR，DSA加速器厂商可以快速的搭建其定制化算子编译器。如Google论文 :cite:`vasilache2022composable`中所述，当前的算子编译器大多提供了一整套自顶向下的编译优化pass，包括调度优化、切分优化、窥孔优化、后端优化、指令生成等，彼此之间大多无法复用，导致新的场景中通常又得从头开发。而在MLIR中，将功能相近的IR优化pass封装为方言（Dialect），并且提供了多个代码生成相关的基础方言，如vector、memref、tensor、scf、affine、linalg等。硬件厂商可以基于这些方言，快速构建一整套lower优化及codegen流程。如 :numref:`MLIR_Lowing`所示，利用scf、affine、linalg等方言，对结构化的计算IR完成循环并行优化、切分、向量化等，最后基于LLVM完成指令映射。

 ![MLIR_Lowing](../img/ch06/MLIR-Lowing.svg)
 :width:`800px`
@@ -47,25 +47,25 @@

 -   **华为TBE/AKG**

- TBE（Tensor Boost Engine）是华为的Ascend芯片及其CANN软件栈基于TVM 开发的一套算子编译优化工具，用于对Ascend芯片进行调度优化、指令映射、及后端pass优化等。不仅提供了一个优化过的神经网络标准算子库，同时还提供了算子开发能力及融合能力。通过TBE提供的API和自定义算子编程开发界面可以完成相应神经网络算子的开发，帮助用户较容易的去使能硬件加速器上的AI_CORE 相关指令，以实现高性能的神经网络计算。为了简化算子开发流程，TBE还实现了一个Auto Schedule工具，开放了自定义算子编程DSL，用于自动完成复杂算子的调度生成。此外，TBE还实现了端到端的动态shape算子编译能力。
+张量加速引擎（Tensor Boost Engine，TBE）是华为的Ascend芯片及其CANN软件栈基于TVM 开发的一套算子编译优化工具，用于对Ascend芯片进行调度优化、指令映射、及后端pass优化等。不仅提供了一个优化过的神经网络标准算子库，同时还提供了算子开发能力及融合能力。通过TBE提供的API和自定义算子编程开发界面可以完成相应神经网络算子的开发，帮助用户较容易的去使能硬件加速器上的AI Core指令，以实现高性能的神经网络计算。为了简化算子开发流程，TBE还实现了一个Auto
+Schedule工具，开放了自定义算子编程DSL，用于自动完成复杂算子的调度生成。此外，TBE还实现了端到端的动态形状算子编译能力。
 
 ![TBE](../img/ch06/TBE.svg)
 :width:`800px`
 :label:`tbe`

- AKG则是MindSpore社区的开源算子编译工具。与上述介绍的算子编译器不同，AKG基于Polyhedral多面体编译技术 :cite:`bastoul2004code`，支持在CPU/GPU/Ascend多硬件上自动生成满足并行性与数据局部性的调度。Polyhedral编译技术的核心思想是将程序中循环的迭代空间映射为高维空间多面体，通过分析语句读写依赖关系，将循环调度优化问题转换为整数规划求解问题。
- AKG的编译流程如 :numref:`akg`所示，主要包含程序规范化、自动调度优化、指令生成、后端优化几个模块。AKG同样基于TVM实现，支持TVM compute/Hybrid DSL编写的算子表示，以及MindSpore图算融合模块优化后的融合子图。通过IR规范化，将DSL/子图IR转换为polyhedral编译的调度树。在polyhedral模块中，利用其提供的调度算法，实现循环的自动融合、自动重排等变换，为融合算子自动生成满足并行性、数据局部性的初始调度。为了能够快速适配不同的硬件后端，我们在poly模块内将优化pass识别为硬件无关的通用优化与硬件相关的特定优化，编译时按照硬件特征拼接组合，实现异构硬件后端的快速适配。
+AKG则是MindSpore社区的开源算子编译工具。与上述介绍的算子编译器不同，AKG基于Polyhedral多面体编译技术 :cite:`bastoul2004code`，支持在CPU、GPU和Ascend多种硬件上自动生成满足并行性与数据局部性的调度。Polyhedral编译技术的核心思想是将程序中循环的迭代空间映射为高维空间多面体，通过分析语句读写依赖关系，将循环调度优化问题转换为整数规划求解问题。 AKG的编译流程如 :numref:`akg`所示，主要包含规范化、自动调度优化、指令映射、后端优化几个模块。AKG同样基于TVM实现，支持TVM compute/Hybrid DSL编写的算子表示，以及MindSpore图算融合模块优化后的融合子图。通过IR规范化，将DSL/子图IR转换为Polyhedral编译的调度树。在Poly模块中，利用其提供的调度算法，实现循环的自动融合、自动重排等变换，为融合算子自动生成满足并行性、数据局部性的初始调度。为了能够快速适配不同的硬件后端，在Poly模块内将优化pass识别为硬件无关的通用优化与硬件相关的特定优化，编译时按照硬件特征拼接组合，实现异构硬件后端的快速适配。

 ![AKG](../img/ch06/akg.png)
 :width:`800px`
 :label:`akg`

- 在polyhedral模块中，实现了算子的自动调度生成、自动切分以及自动数据搬移。为了进一步提升算子的性能，我们针对不同硬件后端开发了相应的优化pass，如Ascend后端中实现数据对齐、指令映射，GPU后端中实现向量化存取，插入同步指令等，最终生成相应平台代码。
+在Poly模块中，实现了算子的自动调度生成、自动切分以及自动数据搬移。为了进一步提升算子的性能，针对不同硬件后端开发了相应的优化pass，如Ascend后端中实现数据对齐、指令映射，GPU后端中实现向量化存取，插入同步指令等，最终生成相应平台代码。

 ### 硬件加速器的多样化编程方法
 :label:`diversified-programming-title`

-矩阵乘法运算作为深度学习网络中占比最大的计算，对其进行优化是十分必要的。因此本节将统一以矩阵乘法$D[M, N] = C[M, N] + A[M, K] \times B[K, N]$为实例，对比介绍如何通过不同编程方式使能加速器。
+矩阵乘法运算作为深度学习网络中占比最大的计算，对其进行优化是十分必要的。因此本节将统一以广义矩阵乘法为实例，对比介绍如何通过不同编程方式使能加速器。广义矩阵乘法指GEMM（General Matrix Multiplication），即$\bm{C} = \alpha \bm{A}\times \bm{B} + \beta \bm{C}$，其中$\bm{A}\in\mathbb{R}^{M\times K}, \bm{B}\in\mathbb{R}^{K\times N}, \bm{C}\in\mathbb{R}^{M\times N}$。

 ![矩阵乘法GEMM运算](../img/ch06/gemm.svg)
 :width:`800px`
@@ -75,8 +75,9 @@

 -   **算子库层级**

-在上述不同层级的编程方式中，直接调用算子加速库使能加速器无疑是最快捷高效的方式。NVIDIA提供了cuBLAS/cuDNN两类算子计算库，cuBLAS提供了使能Tensor Core单元的接口，用以加速矩阵乘法(GEMM)运算，cuDNN提供了对应接口加速卷积(CONV)运算等。
-以 :numref:`accelerator-programable-title`小节的GEMM运算为例，与常规CUDA调用cuBLAS算子库相似，通过cuBLAS加速库使能Tensor Core步骤包括：
+在上述不同层级的编程方式中，直接调用算子加速库使能加速器无疑是最快捷高效的方式。NVIDIA提供了cuBLAS/cuDNN两类算子计算库，cuBLAS提供了使能张量计算核心的接口，用以加速矩阵乘法(GEMM)运算，cuDNN提供了对应接口加速卷积(CONV)运算等。
+以 :numref:`accelerator-programable-title`小节的GEMM运算为例，与常规CUDA调用cuBLAS算子库相似，通过cuBLAS加速库使能张量计算核心步骤包括：
+
 1.  创建cuBLAS对象句柄且设置对应数学计算模式

 ```cpp
@@ -115,14 +116,14 @@ cudaFree(devPtrA);
 cudaDestroy(handle);
 ```

-当然，由于加速器一般会有矩阵形状、数据类型、排布方式等限制，因此在调用句柄和函数接口时要多加注意。如本例中，cuBLAS计算模式必须设置为$CUBLAS\_TENSOR\_OP\_MATH$，步长必须设置为8的倍数，输入数据类型必须为$CUDA\_R\_16F$等。按照如上方式即可通过cuBLAS算子库对 :numref:`accelerator-programable-title`实例使能Tensor Core加速器，通过NVIDIA官方数据可知，该方式对于不同矩阵乘法计算规模，平均有4～10倍的提升，且矩阵规模越大，加速器提升效果越明显。
+当然，由于加速器一般会受到矩阵形状、数据类型、排布方式等限制，因此在调用句柄和函数接口时要多加注意。如本例中，cuBLAS计算模式必须设置为$CUBLAS\_TENSOR\_OP\_MATH$，步长必须设置为8的倍数，输入数据类型必须为$CUDA\_R\_16F$等。按照如上方式即可通过cuBLAS算子库对 :numref:`accelerator-programable-title`实例使能张量计算核心，通过NVIDIA官方数据可知，该方式对于不同矩阵乘法计算规模，平均有4～10倍的提升，且矩阵规模越大，加速器提升效果越明显。

 该方式由于能够隐藏体系结构细节，易用性较好，且一般官方提供的算子库吞吐量较高。但与此同时，这种算子颗粒度的库也存在一些问题，如不足以应对复杂多变的网络模型导致的算子长尾问题（虽然常规形式算子占据绝大多数样本，但仍有源源不断的新增算子，因其出现机会较少，算子库未对其进行有效优化。），以及错失了较多神经网络框架优化（如算子融合）的机会。

 -   **编程原语层级**

-第二种加速器编程方式为编程原语使能加速器，如通过在Device端调用CUDA WMMA (Warp Matrix Multiply Accumulate) API接口。以线程束（即Warp，是调度的基本单位）为操纵对象，使能多个Tensor Core单元。该方式在CUDA 9.0中被公开，程序员可通过添加API头文件的引用和命名空间定义来使用上述API接口。基于软硬件协同设计的基本思想，该层级编程API的设计多与架构绑定，如WMMA操纵的总是$16\times16$大小的矩阵块，并且操作一次跨两个TensorCore进行处理，本质是与TensorCore如何集成进SM中强相关的。针对Float16输入数据类型，NVIDIA官方提供了三种不同矩阵规模的WMMA乘累加计算接口，分别为$16\times16\times16$，$32\times8\times16$，$8\times32\times16$。
-该API接口操纵的基本单位为Fragment，是一种指明了矩阵含义（乘法器/累加器）、矩阵形状（$WMMA\_M, WMMA\_N, WMMA\_K$）、数据类型（Half/ Float）、排布方式（$row\_major/ col\_major$）等信息的模板类型，包括如下：
+第二种加速器编程方式为编程原语使能加速器，如通过在Device端调用CUDA WMMA (Warp Matrix Multiply Accumulate) API接口。以线程束（即{Warp}，是调度的基本单位）为操纵对象，使能多个张量计算核心。该方式在CUDA 9.0中被公开，程序员可通过添加API头文件的引用和命名空间定义来使用上述API接口。基于软硬件协同设计的基本思想，该层级编程API的设计多与架构绑定，如在Volta架构中WMMA操纵的总是$16\times16$大小的矩阵块，并且操作一次跨两张量计算核心进行处理，本质是与张量计算核心如何集成进SM中强相关的。在Volta架构下，针对FP16输入数据类型，NVIDIA官方提供了三种不同矩阵规模的WMMA乘累加计算接口，分别为$16\times16\times16$，$32\times8\times16$，$8\times32\times16$。
+该API接口操纵的基本单位为Fragment，是一种指明了矩阵含义（乘法器/累加器）、矩阵形状（$WMMA\_M, WMMA\_N, WMMA\_K$）、数据类型（FP16/FP32）、排布方式（$row\_major/ col\_major$）等信息的模板类型，包括如下：

 ```cpp
 wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, half, wmma::row_major> a_frag;
@@ -130,21 +131,37 @@ wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> b_
 wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> acc_frag;
 wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
 ```
- 使用时，我们需要将待执行乘法操作矩阵块的数据，作为Fragment，由寄存器加载至TensorCore，在将累加Fragment初始化/清零操作后，通过TensorCore单元执行乘累加运算，最后将运算结果的Fragment存回寄存器或其他内存区域。与上述操作对应的，NVIDIA提供了$wmma.load\_matrix\_sync(), wmma.store\_matrix\_sync()$接口用于将参与计算的子矩阵块写入/载出Fragment片段；$wmma.fill\_fragment()$接口用于初始化对应Fragment的数据；$wmma.mma\_sync()$接口用于对Fragment进行乘累加运算。
+使用时，需要将待执行乘法操作矩阵块的数据加载到寄存器，作为Fragment，在将累加Fragment初始化/清零操作后，通过张量计算核心执行乘累加运算，最后将运算结果的Fragment存回到内存。与上述操作对应的，NVIDIA提供了$wmma.load\_matrix\_sync(), wmma.store\_matrix\_sync()$接口用于将参与计算的子矩阵块写入/载出Fragment片段；$wmma.fill\_fragment()$接口用于初始化对应Fragment的数据；$wmma.mma\_sync()$接口用于对Fragment进行乘累加运算。

 -   **指令层级**

-在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。
+在NVIDIA PTX ISA (Instruction Set Architecture)中提供了另一个编程接口，如Volta架构中的$mma.sync.m8n8k4$指令，它使用$M=8, N=8, K=4$的形状配置执行乘累加操作。该API接口操纵的基本单位为数据元素，除了需要指明矩阵尺寸（即修饰符$.m8n8k4$），还需要指明数据的排布类型（用修饰符$.row$或$.col$）以及输入累加器D、矩阵A、矩阵B及输出累加器C的数据格式（使用修饰符$.f32$或$.f16$等）。如要使用PTX指令集，还需要参考官方文档按照相应的语法规则编写，如代码所示。

-![mma指令之线程与矩阵元素映射关系](../img/ch06/ptx.svg)
-:width:`800px`
-:label:`PTX`
+```cpp
+half_t *a, *b;
+float *C, *D;
+unsigned const* A = reinterpret_cast<unsigned const*>(a);
+unsigned const* B = reinterpret_cast<unsigned const*>(b);

-作为一个更细粒度的指令，mma可以组成更加多样化形状的Warp范围的WMMA API接口，可以控制线程束内线程与数据的映射关系，并允许AI编译器自动/手动显式地管理内存层次结构之间的矩阵分解，因此相比于直接应用NVCUDA::WMMA API具有更好的灵活性。
+asm volatile(
+        "mma.sync.aligned.m8n8k4.row.row.f32.f16.f16.f32 "
+        "{%0,%1,%2,%3,%4,%5,%6,%7}, {%8,%9}, {%10,%11}, "
+        "{%12,%13,%14,%15,%16,%17,%18,%19};\n"
+        : "=f"(D[0]), "=f"(D[1]), "=f"(D[2]), "=f"(D[3]), "=f"(D[4]),
+          "=f"(D[5]), "=f"(D[6]), "=f"(D[7])
+        : "r"(A[0]), "r"(A[1]), "r"(B[0]), "r"(B[1]), "f"(C[0]),
+          "f"(C[1]), "f"(C[2]), "f"(C[3]), "f"(C[4]), "f"(C[5]),
+          "f"(C[6]), "f"(C[7]));
+);
+```
+
+使用时，直接将数据元素作为输入传入（对于FP16的数据元素作为$unsigned$类型传入），与上述操作对应的，NVIDIA提供了$ldmatrix$指令用于从共享内存中加载数据到Fragment。
+
+作为一个更细粒度的指令，mma指令可以组成更加多样化形状的Warp范围的WMMA API接口，可以控制线程束内线程与数据的映射关系，并允许AI编译器自动/手动显式地管理内存层次结构之间的矩阵分解，因此相比于直接应用NVCUDA::WMMA API具有更好的灵活性。

 #### 算子编译器编程使能加速器

-基于算子编译器使能加速器实现矩阵乘的流程则对用户更加友好。以在Ascend中使用TBE为例，用户只需基于python定义矩阵乘的tensor信息（数据类型及形状等），调用对应TBE接口即可。如下所示：
+基于算子编译器使能加速器实现矩阵乘的流程则对用户更加友好。以在Ascend中使用TBE为例，用户只需基于python定义矩阵乘的tensor信息（数据类型及形状等），调用对应TBE接口即可。如代码所示：

 ```python
 a_shape = (1024, 256)
--- a/chapter_accelerator/index.md
+++ b/chapter_accelerator/index.md
@@ -1,8 +1,9 @@
 # 硬件加速器

-上一章节，我们详细讨论了计算图的基本组成，生成和执行等关键设计。当前主流深度学习模型大多基于神经网络实现，无论是训练还是推理，都会产生海量的计算任务，尤其是涉及矩阵乘法这种高计算任务的算子。另一方面，通用处理器芯片如CPU，在执行这类算子时通常耗时较大，难以满足训练/推理任务的需求。因此工业界和学术界都将目光投向特定领域的加速器芯片设计，希望以此来解决算力资源不足的问题。
+上一章节详细讨论了后端的计算图优化、算子选择以及内存分配。当前主流深度学习模型大多基于神经网络实现，无论是训练还是推理，都会产生海量的计算任务，尤其是涉及矩阵乘法这种高计算任务的算子。然而，通用处理器芯片如CPU在执行这类算子时通常耗时较大，难以满足训练和推理任务的需求。因此工业界和学术界都将目光投向特定领域的加速器芯片设计，希望以此来解决算力资源不足的问题。

-本章将会着重介绍加速器的基本组成原理，并且以矩阵乘法为例，介绍在加速器上的编程方式及优化方法。最后，介绍由异构算子组成的异构计算图表达与执行方式。
+
+本章将会着重介绍加速器的基本组成原理，并且以矩阵乘法为例，介绍在加速器上的编程方式及优化方法。

 本章的学习目标包括：

@@ -12,8 +13,6 @@

 -   理解编程API的设计理念

-   理解异构硬件加速的表达与执行
-
 ```toc
 :maxdepth: 2