diff --git "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.1-\346\267\261\345\272\246\347\245\236\347\273\217\347\275\221\347\273\234\347\274\226\350\257\221\345\231\250.md" "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.1-\346\267\261\345\272\246\347\245\236\347\273\217\347\275\221\347\273\234\347\274\226\350\257\221\345\231\250.md"
index 0d938bd0..50bfd17e 100644
--- "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.1-\346\267\261\345\272\246\347\245\236\347\273\217\347\275\221\347\273\234\347\274\226\350\257\221\345\231\250.md"
+++ "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.1-\346\267\261\345\272\246\347\245\236\347\273\217\347\275\221\347\273\234\347\274\226\350\257\221\345\231\250.md"
@@ -76,7 +76,7 @@ class Net(nn.Module):
深度神经网络编译器的后端指最终变化后的代码要执行的设备或神经网络加速器,目前常见的支持深度学习的计算设备有CPU、GPU、FPGA、TPU等其它专用加速器。不同的类型的计算设备往往采用完全不同的芯片架构,从而对应的编程模型和优化也完全不同。如CUDA GPU采用的是多个并行的流式处理器(Streaming Multiprocessor)和共享内存的架构,在GPU上执行的代码需要符合SIMT(Single Instruction Multiple Threads)的计算模型;而CPU一般采用的是多核架构,以及多线程模型(如线程池)来实现高性能的计算任务。更进一步,尽管是相同类型的设备,不同的型号都会有不同的硬件参数,如内存大小、算力、带宽等,这些参数都会极大的影响编译优化过程。
## 5.1.3 中间表达
-从前端语言到后端代码的编译过程,和传统编译器类似,需要经过若干中间表达(Intermediate Representation, IR)。目前在神经网络编译器中较为常用的中间表达主要包括计算图(DAG)和算子表达式等。计算图作为连接深度学习框架和前端语言的主要格式,也是标准化深度学习计算模型的常用格式,如ONNX格式即为一种深度学习模型的标准可交换格式,目前主流框架如TensorFlow和PyTorch的大部分程序都可以被转换或导出成ONNX格式。除此之个,每个DNN编译器也会定义自己的计算图格式,一般这些计算图都可以进行等价互相转换。计算图的节点是算子,边表示张量,所有的节点和边构成一张有向无坏图,节点之间的依赖关系表示每个算子的执行顺序。图5-1-5为一个简单的计算图示例。
+从前端语言到后端代码的编译过程,和传统编译器类似,需要经过若干中间表达(Intermediate Representation, IR)。目前在神经网络编译器中较为常用的中间表达主要包括计算图(DAG)和算子表达式等。计算图作为连接深度学习框架和前端语言的主要格式,也是标准化深度学习计算模型的常用格式,如ONNX格式即为一种深度学习模型的标准可交换格式,目前主流框架如TensorFlow和PyTorch的大部分程序都可以被转换或导出成ONNX格式。除此之外,每个DNN编译器也会定义自己的计算图格式,一般这些计算图都可以进行等价互相转换。计算图的节点是算子,边表示张量,所有的节点和边构成一张有向无环图,节点之间的依赖关系表示每个算子的执行顺序。图5-1-5为一个简单的计算图示例。
图5-1-5. 一个简单的计算图示例
diff --git "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.2-\350\256\241\347\256\227\345\233\276\344\274\230\345\214\226.md" "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.2-\350\256\241\347\256\227\345\233\276\344\274\230\345\214\226.md"
index f521575a..76935f9b 100644
--- "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.2-\350\256\241\347\256\227\345\233\276\344\274\230\345\214\226.md"
+++ "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.2-\350\256\241\347\256\227\345\233\276\344\274\230\345\214\226.md"
@@ -49,19 +49,19 @@
图5-2-3. 公共子表达式消除的示例
## 5.2.3 常数传播
-常数传播(constant propagation)就叫常数折叠(constant folding),也是经典编译优化中的常用优化,其主要方法是通过在编译期计算出也是常数表达式的值,用计算出的值来替换原来的表达式,从而节省运行时的开销。在计算图中,如果一个节点的所有输入张量都是常数张量的话,那么这个节点就可以在编译期计算出输入张量,并替换为一个新的常数张量。如图5-2-4为一个常数传播的示例,其中红色方框内的两个节点都可以被提前计算出来,因皮可以在编译期优化掉。值得注意的是,常数传播需要编译器具有计算的能力,甚至对于一些较大的算子还需要能够在加速硬件上(如GPU)上计算,否则优化的过程就会非常的慢。常数传播的优化在深度学习尤其是模型推理的时候非常有用,因为在推理时,模型中的参数张量全部固定为常数张量,大量计算可以在编译期计算好,极大的化简了推理运算时的计算开销。但时,在深度学习的场景中,常数传播有时候也会带来否优化,如增加计算内存甚至计算时间,一个典型的例子就是一个标量常数张量后面跟一个Broadcast的算子时,如果做了常数传播就会增加内存占用,如果后面是访存密集型的算子的话,也会增加内存压力,从而增加计算时间。
+常数传播(constant propagation)就叫常数折叠(constant folding),也是经典编译优化中的常用优化,其主要方法是通过在编译期计算出也是常数表达式的值,用计算出的值来替换原来的表达式,从而节省运行时的开销。在计算图中,如果一个节点的所有输入张量都是常数张量的话,那么这个节点就可以在编译期计算出输入张量,并替换为一个新的常数张量。如图5-2-4为一个常数传播的示例,其中红色方框内的两个节点都可以被提前计算出来,因此可以在编译期优化掉。值得注意的是,常数传播需要编译器具有计算的能力,甚至对于一些较大的算子还需要能够在加速硬件上(如GPU)上计算,否则优化的过程就会非常的慢。常数传播的优化在深度学习尤其是模型推理的时候非常有用,因为在推理时,模型中的参数张量全部固定为常数张量,大量计算可以在编译期计算好,极大的化简了推理运算时的计算开销。但是,在深度学习的场景中,常数传播有时候也会带来负优化,如增加计算内存甚至计算时间,一个典型的例子就是一个标量常数张量后面跟一个Broadcast的算子时,如果做了常数传播就会增加内存占用,如果后面是访存密集型的算子的话,也会增加内存压力,从而增加计算时间。
图5-2-4. 常数传播的示例
## 5.2.4 矩阵乘自动融合
-矩阵乘在深度学习计算图中被广泛应用,如常见的神经网络的线性层、循环神经网络的单元层、注意力机制层等都有大量的矩阵乘法。在同一个网络里,经常会出现形状相同的矩阵乘法,根据一些矩阵的等价规则,如果把些矩阵乘算子融合成一个大的矩阵乘算子,可以更好的利用到GPU的算力,从而加速模型计算。图5-2-5为其中一种常见的矩阵乘自动融合的示例,其中,如果有两个矩阵乘法共享同一个输入张量(图中方框内左侧),我们就可以自动把另个的两个输入张量拼接成一个大的矩阵乘算子(图中方框内右侧),其计算的结果刚好是原算子计算结果的拼接。利用这种规则,图中最右侧的GRU网络中的两组矩阵乘算子可以分别融合成两个大的矩阵乘算子。类似的融合规则还有BatchMatMul,可以把两个相同形状的矩阵拼接成一个新的BatchMatMul算子。
+矩阵乘在深度学习计算图中被广泛应用,如常见的神经网络的线性层、循环神经网络的单元层、注意力机制层等都有大量的矩阵乘法。在同一个网络里,经常会出现形状相同的矩阵乘法,根据一些矩阵的等价规则,如果把些矩阵乘算子融合成一个大的矩阵乘算子,可以更好的利用到GPU的算力,从而加速模型计算。图5-2-5为其中一种常见的矩阵乘自动融合的示例,其中,如果有两个矩阵乘法共享同一个输入张量(图中方框内左侧),我们就可以自动把另外的两个输入张量拼接成一个大的矩阵乘算子(图中方框内右侧),其计算的结果刚好是原算子计算结果的拼接。利用这种规则,图中最右侧的GRU网络中的两组矩阵乘算子可以分别融合成两个大的矩阵乘算子。类似的融合规则还有BatchMatMul,可以把两个相同形状的矩阵拼接成一个新的BatchMatMul算子。
图5-2-5. 矩阵乘自动融合的示例
## 5.2.5 算子融合
-上小节中介绍的算子融合方法是针对矩阵乘算子特有的,在深度学习模型中,针对大量的小算子的融合都可以提高GPU的利用率,减少内核启动开销、减少访存开销等好处。例如,Element-wise的算子(如Add,Mul,Sigmoid,Relu等)其计算量非常小,主要计算瓶颈都在内存的读取和写出上,如果前后的算子能够融合起来,前面算子的计算结果就可以直接被后面算子在寄存器中使用,避免数据在内存的读写,从而提交整体计算效率。图5-2-6展示了一个Mul算子和一个Add算子融合的示例,图5-2-7为其对应的融合前后的CUDA代码示例,在没有融合前,执行两个算子需要启动两个GPU内核,前一个计算的结果需要写出到主存中,下一个内核计算的时候需要再次读取的计算核上。然后,融合后的代码只需要启动一个内核,并且可以有效复用中间计算结果。
+上小节中介绍的算子融合方法是针对矩阵乘算子特有的,在深度学习模型中,针对大量的小算子的融合都可以提高GPU的利用率,减少内核启动开销、减少访存开销等好处。例如,Element-wise的算子(如Add,Mul,Sigmoid,Relu等)其计算量非常小,主要计算瓶颈都在内存的读取和写出上,如果前后的算子能够融合起来,前面算子的计算结果就可以直接被后面算子在寄存器中使用,避免数据在内存的读写,从而提高整体计算效率。图5-2-6展示了一个Mul算子和一个Add算子融合的示例,图5-2-7为其对应的融合前后的CUDA代码示例,在没有融合前,执行两个算子需要启动两个GPU内核,前一个计算的结果需要写出到主存中,下一个内核计算的时候需要再次读取到计算核上。然后,融合后的代码只需要启动一个内核,并且可以有效复用中间计算结果。
图5-2-6. 算子融合的示例
diff --git "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.3-\345\206\205\345\255\230\344\274\230\345\214\226.md" "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.3-\345\206\205\345\255\230\344\274\230\345\214\226.md"
index e9e1a278..2ea1f71f 100644
--- "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.3-\345\206\205\345\255\230\344\274\230\345\214\226.md"
+++ "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.3-\345\206\205\345\255\230\344\274\230\345\214\226.md"
@@ -30,12 +30,12 @@
上面的方法中只考虑了张量放置在加速器(如GPU)的内存中,而实际上如果内存不够的话,我们还可以将一部分张量放置到外存中(如CPU的内存中),等需要的时候再移动回GPU的内存中即可。虽然从CPU的内存到GPU的内存的拷贝延时和带宽都比较受限,但是因为计算图中有些张量的产生到消费中间会经过较长的时间,我们可以合理安排内存的搬运时机使得其和其它算子的计算重叠起来。
-图5-3-2. 利用整数线性规划优化计算图内存分配的示
+图5-3-2. 利用整数线性规划优化计算图内存分配的示例
-给定上述假设以及必要的数据(如每个内核的执行时间、算子的执行顺序等),关于每个张量在什么时间放在什么地方的问题就可以被形式化的描述成一个最优化问题。AutoTM[]就提出了一种把计算图中张量内在异构内存环境中规的问题划建模成一个整数线性规划的问题并进行求解。图5-3-2展示了一个利用整数线性规划优化计算图内存分配的优化空间示例,图中每一行表示一个张量,每一列表示算子的执行顺序。每一行中,黄色Source表示张量的生成时间,紫色的SINK表示张量被消费的时间,每个张量都可以选择是在内存中(DRAM)还是外存(PMM)中。那么问题优化目标为就是给定任意的计算图最小化其执行时间,约束为主存的占用空间,优化变量就是决定放在哪个存储中,在有限的节点规模下,这个问题可以通过整数线性规划模型求解。同时,该文章中还扩展该方法并考虑了更复杂的换入换出的情形。
+给定上述假设以及必要的数据(如每个内核的执行时间、算子的执行顺序等),关于每个张量在什么时间放在什么地方的问题就可以被形式化的描述成一个最优化问题。AutoTM[]就提出了一种把计算图中张量在异构内存环境中规划的问题建模成一个整数线性规划的问题并进行求解。图5-3-2展示了一个利用整数线性规划优化计算图内存分配的优化空间示例,图中每一行表示一个张量,每一列表示算子的执行顺序。每一行中,黄色Source表示张量的生成时间,紫色的SINK表示张量被消费的时间,每个张量都可以选择是在内存中(DRAM)还是外存(PMM)中。那么问题优化目标为就是给定任意的计算图最小化其执行时间,约束为主存的占用空间,优化变量就是决定放在哪个存储中,在有限的节点规模下,这个问题可以通过整数线性规划模型求解。同时,该文章中还扩展该方法并考虑了更复杂的换入换出的情形。
## 5.3.3 张量重计算
-深度学习计算图的大多算子都是确定性的,即给定相同的输入其计算结果也是相同的。因此,我们可以进一步利用这个特点来优化内存的使用。当我们对连续的多个张量决定换入换出的方案时,如果产生这些张量的算子都具有计算确定性的话,我们可以选择只换出其中一个或一少部分张量,并把剩下的张量直接释放,当到了这些张量使用的时机,我们可以再换入这些少量的张量,并利用确定性的特点重新计算之前被释放的张量,这样就可以一定程序上缓解CPU和GPU之前的贷款压力,也为内存优化提供了更大的空间。如果考虑上换入换出,内存优化方案需要更加仔细的考虑每个算子的执行时间,从而保证重计算出的张量在需要的时候能及时的计算完成。
+深度学习计算图的大多算子都是确定性的,即给定相同的输入其计算结果也是相同的。因此,我们可以进一步利用这个特点来优化内存的使用。当我们对连续的多个张量决定换入换出的方案时,如果产生这些张量的算子都具有计算确定性的话,我们可以选择只换出其中一个或一部分张量,并把剩下的张量直接释放,当到了这些张量使用的时机,我们可以再换入这些少量的张量,并利用确定性的特点重新计算之前被释放的张量,这样就可以一定程序上缓解CPU和GPU之前的带宽压力,也为内存优化提供了更大的空间。如果考虑上换入换出,内存优化方案需要更加仔细的考虑每个算子的执行时间,从而保证重计算出的张量在需要的时候能及时的计算完成。
## 小结与讨论
diff --git "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.4-\345\206\205\346\240\270\344\274\230\345\214\226.md" "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.4-\345\206\205\346\240\270\344\274\230\345\214\226.md"
index c648f726..4e0bbabb 100644
--- "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.4-\345\206\205\346\240\270\344\274\230\345\214\226.md"
+++ "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.4-\345\206\205\346\240\270\344\274\230\345\214\226.md"
@@ -52,7 +52,7 @@ for (int i= 0; i < n; ++i)
```
图5-4-2. 一个张量加算子的调度示例
-可以看到,上面生成的内核代码只是一个简单的循环,实际中这样的代码往往性能不好。我们希望对上述循环进行一系列的变化,如把一个循环拆分成两重循环、或者把两个循环合并一个循环、或者把两个循环的顺序颠倒等等。为了方便这些优化,算子编译器也提供了一些相应的调度操作接口,如下图中的split操作即可以上述循环按照32为因子进行拆分成内个两重循环,如图5-4-3所示。
+可以看到,上面生成的内核代码只是一个简单的循环,实际中这样的代码往往性能不好。我们希望对上述循环进行一系列的变化,如把一个循环拆分成两重循环、或者把两个循环合并一个循环、或者把两个循环的顺序颠倒等等。为了方便这些优化,算子编译器也提供了一些相应的调度操作接口,如下图中的split操作即可以上述循环按照32为因子进行拆分成一个两重循环,如图5-4-3所示。
```
# 在TVM中创建一个默认调度的示例
@@ -105,7 +105,7 @@ if (i < n)
## 5.4.3 自动调度搜索与代码生成
-有了算子表达式和对表达式的调度机制,我们就可以较容易的在一个新的硬件设备上生成一个算子的内核代码了。然而,我们可以看到,在调度的时候,有非常多种决定需要抉择,而且这些决定都会根据硬件的不同而产生不一样的性能影响,这些都需要经验非常丰富的专家才能知道一个较好的调度方案。为了进一步克复这个问题,一类利用机器学习进行自动调度搜索的方法被广泛应用。
+有了算子表达式和对表达式的调度机制,我们就可以较容易的在一个新的硬件设备上生成一个算子的内核代码了。然而,我们可以看到,在调度的时候,有非常多种决定需要抉择,而且这些决定都会根据硬件的不同而产生不一样的性能影响,这些都需要经验非常丰富的专家才能知道一个较好的调度方案。为了进一步克服这个问题,一类利用机器学习进行自动调度搜索的方法被广泛应用。
图5-4-5. 自动调度搜索与代码生成
@@ -136,4 +136,4 @@ if (i < n)
8. Ansor: Generating high-performance tensor programs for deep learning.
-9. Flextensor: An automatic schedule exploration and optimization framework for tensor computation on heterogeneous system.
\ No newline at end of file
+9. Flextensor: An automatic schedule exploration and optimization framework for tensor computation on heterogeneous system.
diff --git "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.5-\347\256\227\345\255\220\350\260\203\345\272\246\344\274\230\345\214\226.md" "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.5-\347\256\227\345\255\220\350\260\203\345\272\246\344\274\230\345\214\226.md"
index 83c989fe..a8a99993 100644
--- "a/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.5-\347\256\227\345\255\220\350\260\203\345\272\246\344\274\230\345\214\226.md"
+++ "b/Textbook/\347\254\2545\347\253\240-\346\267\261\345\272\246\345\255\246\344\271\240\346\241\206\346\236\266\347\232\204\347\274\226\350\257\221\344\270\216\344\274\230\345\214\226/5.5-\347\256\227\345\255\220\350\260\203\345\272\246\344\274\230\345\214\226.md"
@@ -42,7 +42,7 @@ __device__ float MatMul(float *a, float *b, float *c, int m, int n, int k) {
}
```
-为了按照前面讲到的的方法进行这三个算子的融合,我们需要将上述三个函数生成到同一个全局核函数内,如下图示例。值的注意的是,为了保证任意算子之前有正确的数据依赖,我们有时候需要在两个算子之间插入一个全局的数据同步。
+为了按照前面讲到的的方法进行这三个算子的融合,我们需要将上述三个函数生成到同一个全局核函数内,如下图示例。值得注意的是,为了保证任意算子之前有正确的数据依赖,我们有时候需要在两个算子之间插入一个全局的数据同步。
```
//融合三个算子后的全局内核函数
__global__
@@ -80,7 +80,7 @@ void kernel_0(float *A, float *B, float *C) {
为了更好的解决这种问题,我们就是需要一种能根据计算流图中的并行度以及算子内部的并行度的整体信息来进行一个全局的任务调度方法。本章中以Rammer的技术为例,介绍一种全局算子调度的优化方法。
首先,在计算表达层,为了打开现有算子的黑盒实现,Rammer引入rOperator来代替原有算子抽象,暴露出每一个算子中的所有并行任务(rTask)。
-在硬件层,引入虚拟设备(vDevice)的抽象,并提供计算单元(vEU)级别的的粒度调度接口,中以允许将一个rTask调度到任意指定的vEU上。然而,rTask粒度的调度可能带来更严重的调度开销,Rammer利用DNN计算性能有较强的确定性,即算子的计算时间在数据流图运行前就可以通过测量得到。因此,在编译时可以将整个数据流图的所有rTask静态的编排成一个确定性执行方案,通过vDevice映射到物理Device的执行单元进行执行,如图5-5-2所示。
+在硬件层,引入虚拟设备(vDevice)的抽象,并提供计算单元(vEU)级别的的粒度调度接口,可以允许将一个rTask调度到任意指定的vEU上。然而,rTask粒度的调度可能带来更严重的调度开销,Rammer利用DNN计算性能有较强的确定性,即算子的计算时间在数据流图运行前就可以通过测量得到。因此,在编译时可以将整个数据流图的所有rTask静态的编排成一个确定性执行方案,通过vDevice映射到物理Device的执行单元进行执行,如图5-5-2所示。
图5-5-2. 全局算子调度编译