背景

要让AI芯片支持深度学习架构，要将深度学习架构等部署到芯片上就需要将深度学习架构中的这些代码编译成芯片支持的指令集，所以要从头到尾设计一套软件栈，做一套全栈的优化。所以现在的许多深度学习的架构只能在某一些厂商的GPU设备上获得加速，这种支持依赖于特定的GPU库，当未来越来越多的加速器出现时，硬件设备的厂商对于深度学习架构的支持便会变得越来越困难。

原来的框架：PyTorch、Tensorflow运用于GPU加速，存在很大的局限，部署运用其他平台需要很多手动的工作

CPU、GPU和类似于处理器的加速器需要不同的片上存储器架构和计算原语。在生成优化代码时必须解决这个分歧。
基本优化的挑战：深度学习用于深度学习的优化编译器需要同时提供高级和低级的优化
总结了graph level and tensor operator level 的主要的四大挑战
① 高级数据流重写: 不同的硬件设备不同的内存层次结构，因此启用融合操作符和优化数据布局的策略对于优化内存访问至关重要
② 跨线程的内存重用：现代GPU和专门的加速器都有可以共享计算的内存。传统的无共享嵌套并行模型已不再是最优的。优化内核需要线程之间在共享内存上的协作。
③ 张量化的计算函数：最新的硬件提供了矢量操作之外的新指令，因此调度程序必须将计算分解为张量计算在内的函数，而不是标量或者矢量代码
④ 延迟隐藏：传统的同步多线程和自动管理的缓存结构隐式地隐藏了现代cpu / gpu的延迟，而专门的加速器设计通常倾向于更精简的控制，并将大部分调度复杂性丢给编译器堆栈。不过，调度必须谨慎地执行，以隐藏内存访问延迟。

TVM

TVM一种端到端优化堆栈，目标是很容易的将深度学习架构部署到不同内存结构、计算单元的所有的硬件设备中（包括GPU，FPGA 和 ASIC（如谷歌 TPU），也包括嵌入式设备）

目的

TVM旨在架起多种深度学习前端框架与后端硬件之间的桥梁，能够将深度学习模型高效地部署到异构的后端上并进行自动调优。轻松地将深度学习工作负载部署到各种硬件目标上，包括嵌入式设备、gpu、fpga和asic(例如TPU)，它们在内存组织、计算等方面存在显著差异
TVM效果：跨硬件后端提供的性能可以与低功耗CPU和服务器类gpu的最新库相媲美

TVM提出了两个优化层:一个是计算图优化层，用于解决第一个调度难题;另一个是具有新的调度原语的张量优化层，用于解决剩下的三个难题
通过结合这些优化层，TVM可以从大多数深度学习框架中获取模型描述，执行高级和低级的联合优化，并为后端生成特定于硬件的优化代码

优化计算图（Optimizing Computational Graphs）

TVM利用一种计算图表示来处理高层次的优化:节点表示对静态已知维数的张量的操作，节点表示对静态已知维数的张量的操作，而边表示张量操作之间的数据流数据依赖性。

① 算子融合：

将多个操作符融合在一起是一种优化，可以大大减少执行时间，特别是在gpu和专门的加速器。其思想是将多个操作符组合到一个内核中，而不将中间结果保存到全局内存中。用cache提高cache命中率，速度得到提升，另一方面，可以做些计算优化放到某些算子中间，提高本地性

下图有个算子融合里图例

② 数据布局转换：

张量运算是计算图的基本运算符。计算中涉及的张量在不同的操作中可能有不同的布局要求，不同的硬件设备可能具有截然不同的内存层次结构
因此，融合算子与优化数据布局的策略对于优化内存访问至关重要。在目标硬件平台中，给定内存层次结构的约束，为每个操作符指定首选的数据布局，若数据布局不匹配，改变计算图中张量的数据布局方式以适应目标硬件设备并提升执行效率。
下图提供一个

图级优化的局限

虽然高层次的数据流图优化可以极大地提高深度学习工作负载的效率，但它们的有效性仅与操作库提供的一样。目前，支持操作员融合的深度学习框架很少，它们需要操作员库提供融合模式的实现。随着越来越多的网络运营商的引入，可能融合的内核数量会急剧增加。当以数量不断增加的硬件后端为目标时，这种方法不再是可持续的，因为融合模式实现的所需数量与需要支持的数据布局、数据类型和硬件加速特性的数量结合在一起增长。对于这种庞大的后端特定操作符空间，手工操作符内核是不可行的。

张量优化操作 (OptimizingTensor Operations)

TVM如何为各种硬件后端生成同一操作符的微调版本

① 张量表达式语言：

我们引入了一种数据张量表达语言来支持自动代码生成。
Tensor Expression 支持常见的算术运算和数学运算，并涵盖了常见的算子。该语言没有指定循环结构和许多其他执行细节，并且它提供了为各种后端添加硬件感知优化的灵活性。根据解耦计算和调度的原则，用调度表示从张量表达式到底层代码的映射。
（就是给定一个张量没有指明是如何实现这个操作的步骤）
与高级计算图语言不同，在高级计算图语言中张量运算的实现是不透明的，每个运算都用索引公式表达式语言描述，如图所示

    TVM中用于算法描述的张量表达式语言。每个计算操作包含输出张量的形状，以及描述如何计算输出张量的每个元素的表达式

我们的张量表达式语言支持常见的算术和数学运算，这些运算可以在常见的语言(如c)中找到。我们显式地引入了交换约简运算符，以便在多线程中轻松地安排交换约简
我们进一步介绍了一个高阶扫描运算符，它可以结合基本计算运算符来形成随时间变化的循环计算
TVM计算操作还支持张量元组之间的简化，这使得支持像argmax这样的函数变得很容易。这种表示可以描述高级数据流图中使用的所有张量操作，并涵盖了深度学习中常见的模式。

② 调度空间

					应用于CPU、GPU和深度学习加速器设计的典型优化

每个优化的低级程序都是不同调度策略组合的结果，给内核编写器带来了很大的负担。
调度是将计算描述降低到后端优化实现的特定规则。其思想是对调度空间和用于遍历该空间的转换进行正确建模，从而提供生成低级代码的不同方法。TVM的调度空间如图所示
为了快速探索计划空间，我们需要提供有效的计划原语来转换计划

// 例如一段最原始的TVM代码：c = tvm.compute((n,), lambda i:A[i]+B[i])s = tvm.create_schedule(c.op)// 生成得到的代码是：for(int i=0; i<n; ++i){c[i] = A[i] + B[i]
} 
###########################
###########################
// 加上额外的调度c = tvm.compute((n,), lambda i:A[i]+B[i])s = tvm.create_schedule(c.op)xo,xi = s[c].split(s[c].axis[0],factor=32)// 再生成的代码就变成了
for(int xo=0; xo<ceil(n/32); ++xo){for(int xi=0; xi<32; ++xi){int i = xo * 32 + xi;if(i<n){c[i] = A[i] + B[i];}}
}###########################
###########################// 甚至于还可以绑定中间的xo和xi到特定的变量上：c = tvm.compute((n,), lambda i:A[i]+B[i])s = tvm.create_schedule(c.op)xo,xi = s[c].split(s[c].axis[0],factor=32)s[c].recoder(xi,xo)s[c].bind(xo,tvm.thread_axis("blockIdx.x"))s[c].bind(xi.tvm.thread_axis("threadIdx.x"))

compute决定了一段程序的输入输出的关系，而schedule决定了计算的顺序等因素。
schedule虽然多种多样，也可以自由组合，但是无论如何也不会改变计算的结果。

在TVM目前（2019.10.23）提供的schedule中，主要常用的可以归纳为如下两种：

I : 针对循环的变换：

TVM提供了循环展开(unroll)，循环交换顺序(reorder)我们常见的循环优化，以及bind（上文提到过）等schedule原语。除此之外，为了更好的发挥他们的威力，也提供了循环切割(split),循环合并(fuse)，循环切块(tiling)等原语

II : 访存调整：

程序执行的快慢会与cache命中率有关

III ：提出新的调度解决GPU和专门的硬件加速器的问题

③ 嵌套并行与协作

并行编程是提高深度学习工作中计算密集型核效率的关键。现代GPU提供了大量的并行性，这要求我们将并行编程模型放入调度转换中。大多数现有的解决方案采用被称为嵌套并行程序的并行编程模型，这是fork-join并行的一种形式。具体来说，我们可以使用parallel schedule原语来并行化数据并行任务。每个并行任务可以进一步递归地细分为子任务，以利用目标体系结构上的多级线程层次结构
我们称这个模型为无共享嵌套并行，因为一个工作线程不能在同一并行计算阶段查看它同级线程的数据。同级线程之间的交互发生在连接阶段，此时子任务已经完成，下一阶段可以使用前一阶段产生的数据。这种编程模型不允许线程之间相互协作，以便在同一并行阶段内执行集体任务
在GPU上的矩阵乘法要求将工作划分为块，这些块分布在线程组中，然后分配给每个单独的线程。在无共享嵌套并行性下，每个线程在减少阶段必须独立地获取计算平摊所需的数据。
不共享方法的一个更好的替代方法是跨线程协作地获取数据。这种模式在使用CUDA、OpenCL和Metal等语言的GPU编程中很常见，但还没有实现到schedule原语中。我们将内存作用域的概念引入到调度空间中，以便可以将一个阶段标记为共享。如果没有内存作用域，自动作用域推断将把相关阶段标记为线程本地。共享任务需要计算组中所有工作线程的依赖关系。通过使用线程绑定原语将加载任务分配到同一组线程中，可以有效地调度输入数据的加载。
因此线程可以跨加载和计算阶段持久存在。这个改进的实现需要额外的编译器支持。具体来说，绑定推理算法需要能够通过合并所有合作线程的任务来推断共享任务的边界。此外，需要适当地插入内存同步屏障，以确保共享加载的数据对使用者可见

	引入内存作用域以支持并行作业之间的任务共享由此产生的新调度可以更好地利用现代gpu中的共享内存层次结构。

④ 张量化：概括硬件接口

最新的硬件带来了超越向量运算的新指令集，调度必须使用这些指令才能从加速中获益。张量化和向量化存在区别，张量计算基元的输入是多维的，具有固定的或可变的长度，并指示不同的数据布局，更重要的是，我们不能求助于一套固定的原语，因为新的深度学习加速器正以它们自己的张量指令出现。因此，我们需要一个解决方案，作为未来支持新一代专业加速器。
为了解决这个挑战，我们将硬件接口与调度分开。特别地，我们引入了一个张量本征声明机制。我们可以使用张量表达式语言来声明每个新硬件的内在行为，以及与之相关的降低规则。此外，我们引入了一个张量调度原语来代替计算单位与相应的张量本质。编译器用硬件声明匹配计算模式，并将其降低为对应的硬件内部声明。
张量表达式语言既描述了用户预期的计算描述，也描述了硬件所暴露的抽象。紧张调度生成的代码与高性能计算中的常见做法一致:将复杂的操作分解为重复的微内核调用序列。因此，我们还可以使用tensorize原语来利用手工组装的微内核，这在某些平台上是有益的。
张量化：不光利用张量表达式表示高层的运算，同样利用同样的张量表达式来表示硬件底层本身的特性。张量化的过程，同时揭露高层张量表达式和底层硬件的特性表述来把高层的表示来映射到底层硬件可以支持的指令上去，这就可以使得TVM可以支持CPU,GPU之类的东西

⑤ 编译器支持延迟隐藏

延迟隐藏是通过重叠内存计算操作去使内存利用率和计算效率最大化。针对不同的硬件终端它需要不同的策略。在CPU中通过同步多线程来实现延迟隐藏，在GPU中通过快速的切换线程去最大化利用功能单元，在某些深度学习加速器中经常将延迟隐藏转移到堆栈编译器中去完成。在硬件设备上实现延迟隐藏的一种方式是直接控制硬件任务的执行，就是当一个任务完成时另一个独立的任务可以进行的时候，便会在pipeline stages中同步指令或者重写数据。但是这种与硬件直接交互的形式的开销是很大的。为了减少这种开销，TVM提供了一种虚拟的线程调度形式，编程者可以指定一个数据并行程序，然后TVM可以自动的转换成底层的数据依赖程序

对图和调度进行优化后，剩下的任务就是产生能在目标平台以及能够部署在硬件设备上的代码了。

部署问题

1. Code Generation：TVM通过对前端框架的硬件如所用到的CPU，GPU进行分析，将高级语言转换成硬件设备可以支持的C语言。
   对于一个特定的多元数据流声明，坐标系相关的超图和调度树，我们可以通过迭代遍历调度树的方案生成Lowered代码，并且推断出输入Tensor的依赖范围(使用坐标系相关的超图)，接下来生成循环嵌套的low-level代码。low-level代码是通过类C的循环程序具象化而来，在这个过程中使用了一个Halide循环程序数据结构的变体，我们也使用了Halide常用的lowering primitives，像是storage flattening，循环展开，对于GPU和特定深度学习加速器而言，则使用了同步点检测、虚拟线程注入、模块化生成机制。最终，循环程序被翻译成LLVM/CUDA/Metal/OpenCL源代码
2. Runtime Support For GPU programs: TVM会独立的构建host端和device端的模块并且提供一个运行时模块系统来启动Kernel。代码生成算法然后将加速器程序翻译成一系列Runtime API。
3. Autotuning：TVM提供一套新颖的优化框架，这套框架能够为深度学习系统编译高性能底层实现，使得像高维卷积、矩阵乘法和深度卷积这样的复杂操作可以被一套调度器模板进行自动调优。
4. Remote Deployment Profiling：为了嵌入式设备，TVM设计了一套方便性能分析和自动调优的基础设施。传统情况下，嵌入式开发一般是在主机上进行交叉编译然后复制可执行文件到目标设备上执行，上述编译、运行和性能分析的工作都需要手动进行。在编译器堆栈中我们提供了一个远程程序调用：通过RPC接口，我们能够完全在host端完成上面的所有步骤，这样的方式可以极大地加快在嵌入式设备和基于FPGA的加速器上的优化工作。

2020_9_4 第一周组会