1.TVM 自底向上（三）：TE 的码阅概念和编译原理
2.[深入分析CUTLASS系列] 0x02 cutlass 源码分析(一) --- block swizzle 和 tile iterator (附tvm等价code)
3.TVM适配NN编译Compiler缺陷
4.TVM学习－前言篇
5.TVM源语-Compute篇
6.TinyML-TVM是如何驯服Tiny的（上）

TVM 自底向上（三）：TE 的概念和编译原理

       在探讨 TVM 自底向上系列文章的第三篇章节中，我们深入探讨了 Tensor Expression (TE) 的码阅概念及其编译原理。前文已介绍了 TVM 中最接近目标硬件源代码的码阅 IR 表示，即 Tensor IR (TIR)。码阅本文将聚焦 TE，码阅它是码阅c 传奇源码下载位于 Relay IR/TOPI 和 TIR 之间的抽象层次。

       TE 的码阅核心作用之一是提供除 TVMScript 外的另一种方式来构建 TIR AST。TE 通过抽象程度更高的码阅表达式，允许用户以更灵活的码阅方式编写计算逻辑，但 TE 本身无法直接编译为硬件源代码。码阅相反，码阅它需要经过一个称为“lowering”的码阅过程，将 TE 表达式转换为 TIR 的码阅 Primitive Function，之后 TIR 的码阅编译流程得以进行。

       TE 在 TVM 编译流程中的码阅位置如下图所示。首先，使用 TE 来构建计算图，通过调用 `te.create_prim_func` 函数将 TE 表达式转换为 TIR 的 `PrimFunc`。然后，这些 `PrimFunc` 被嵌入到一个新的 `IRModule` 中，该 `IRModule` 再被编译为目标硬件代码。

       以官方文档 Blitz Course to TensorIR 提供的示例为例，使用 TE 实现了一个与前文使用 TVMScript 实现相同的 TIR AST。TE 方法通过 `te.Tensor` 对象构建计算图，随后调用 `te.create_prim_func` 将这些对象转换为 TIR 的 `PrimFunc`。最终，将这些 `PrimFunc` 嵌入到 `IRModule` 中，该模块可被编译为目标硬件执行代码。

       在 TE 中，`tvm.te.Tensor` 是计算图中的数据块，类似于神经网络中的 feature map。TE 提供了两种主要的 `Tensor` 类型：`tvm.te.placeholder` 和 `tvm.te.compute`。`placeholder` 用于计算图的输入节点，而 `compute` 则用于定义计算 Tensor 的操作，基于传入的 lambda 表达式计算结果。

       `PrimExpr` 是 TE 中用于表示 AST 的概念，它是通过 lambda 表达式将 Tensor 和运算转换而成。`PrimExpr` 支持各种数学运算符，并且其内部结构通过 `tir.expr.ProducerLoad` 对象实现，使得 `PrimExpr` 能够自然地表示 AST。通过 `te.compute` 调用将 `PrimExpr` 实例化为 `tvm.te.Tensor`，进一步构建计算图。

       `Operation` 是 TE 中的抽象类，它封装了特定的计算逻辑，例如 `te.compute` 实例化时所对应的计算操作。通过 `Operation`，TE 实现了对计算图中操作的精细控制，包括输入和输出 Tensor 的管理。

       TE 的另一个关键功能是 `te.create_prim_func`，该函数将 TE 表达式转换为 TIR 的 `PrimFunc`。这一过程包括构建一个 Graph，其中包含了所有 Operation 的输出和输入 Tensor 的信息，从而实现 TE 到 TIR 的转换。

       综上所述，TE 作为 TVM 中的高级抽象层，提供了更灵活的web源码开发计算图构建方式，通过与 TIR 的结合实现了高效的目标硬件编译。TE 的引入不仅丰富了 TVM 的编程模型，还为实现不同硬件加速策略提供了更多可能。

[深入分析CUTLASS系列] 0x cutlass 源码分析(一) --- block swizzle 和 tile iterator (附tvm等价code)

       深入探讨CUTLASS系列之block swizzle和tile iterator

       本文聚焦于block swizzle和tile iterator在CUTLASS中的作用。

       block swizzle通过一定的步长进行换行操作，其核心逻辑为取余操作。关注的关键文件包括cutlass/gemm/threadblock/threadblock_swizzle.h和cutlass/gemm/kernel/gemm.h。在GPU中，block的发射顺序为x->y->z，通过位运算实现取余操作，相比直接取余，位运算在开销上更小。

       block swizzle的逻辑分析展示了其在计算过程中的作用，以一个 x x的矩阵乘法为例，不进行block swizzle时，线程块按照n和m轴发射，导致在读取右矩阵的global位置时存在差异，从而影响访存量。进行block swizzle后，单个tile的访存量变小，减少cache miss，提高性能。

       tvm等价代码示例展示了block swizzle的实现方式，简洁明了。

       tile iterator解决的问题在于提供左右矩阵的load/store方法。以conv2d的iterator为例，分析了如何在focus于某一分块时确定每个线程需要被load的位置。重点关注的文件包括cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_analytic.h、cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_optimized.h和cutlass/conv/threadblock/conv2d_tile_iterator.h。分析了shared memory的load过程，以及在不同iterator中的优化方法。

       tile iterator的逻辑分析详细介绍了shared memory的load过程，包括warp的划分、最大访存指令的限制和kStride参数。进一步讨论了analytic和optimized iterator的实现差异，以及如何通过位运算减少scalar操作，提高性能。

       本文总结了block swizzle和tile iterator在CUTLASS中的作用和优化方法，提供了深入理解的途径。希望对相关领域感兴趣的研究者和开发者有所启发。

TVM适配NN编译Compiler缺陷

       在深度学习编译器领域，TVM被设计为缩小深度学习框架与硬件后端性能差距的关键工具。然而，针对自定义神经网络加速器（如VTA）的编译过程中，TVM存在一些缺陷和局限性。

       首先，TVM在处理VTA的自定义架构时，其编译流程存在局限。尽管TVM定义了VTA的指令集和体系结构，但其灵活性较差，比如SRAM配置的固定性导致对硬件资源的适应性不足，特别是对于不同大小的计算阵列和资源分配，编译器未能灵活调整。

       此外，TVM对网络支持有限，10001100变成源码尽管宣称支持主流框架，但在实际应用中，如resnet_v1等模型的编译和运行较为顺利，但扩展到其他模型时可能存在问题，如量化支持的局限和特定bug，这限制了其在更多网络模型上的适用性。

       针对这些缺陷，TVM源码中的静态调度搜索算法成为关键。原生的AutoTVM在线搜索方式在VTA架构固定的情况下，效率不高且不适用于芯片设计阶段。本文作者创新性地实现了静态调度搜索算法，能在编译阶段找到针对特定VTA配置的最优调度，显著提高了性能，且避免了昂贵的FPGA迭代过程，大大缩短了编译时间。

       总结来说，TVM在VTA编译器上需要改进的地方包括：增强SRAM配置的灵活性、优化计算阵列配置的动态调整、扩大网络模型支持范围，以及引入更高效的静态调度策略。作者的成果有望推动TVM更好地适应和优化自定义深度学习硬件，提高编译效率和性能。

TVM学习－前言篇

       本文分享了学习TVM（开源的深度学习模型编译器）的经验与思考。TVM的独特之处在于，它是一个专注于模型部署与推理的引擎，相较于TensorFlow、PyTorch等框架，TVM在行业中的应用日益广泛，许多大型公司都在内部使用TVM。那么，学习TVM有哪些原因？为什么它对学习者来说具有挑战性？又如何有效地学习TVM？

       TVM的学习门槛相对较高。一方面，它假设使用者具备广泛的领域知识，包括AI模型、深度学习框架、Python/C/C++语言以及Linux下的编程知识（如GCC、LLVM、CMake等）。另一方面，TVM在硬件知识上也有一定的要求，如CUDA编程。若在这些领域的知识上有所欠缺，学习TVM时可能会遇到不同程度的困难。

       TVM的调试难度也是一个挑战。作为一个Python与C++混合实现的软件，TVM的调试方式较为复杂。Python层的接口相对容易调试，但C++层的功能实现大多封装在内部，这使得在C++层的调试变得较为困难，往往需要依赖于日志打印或栈回溯等方法。

       TVM概念丰富且抽象，涉及到AST（抽象语法树）概念、IR模型（如Relay IR、TOPI、TE、弹簧软件源码TIR等），对于初学者来说，这些概念可能令人感到困惑。

       学习TVM的过程并不是一帆风顺的。初学者可能会从使用入门、阅读官方文档、尝试各种示例代码入手，逐渐积累经验。但真正的突破往往来自于深入理解源代码，包括阅读、调试、修改源码。这一过程虽然充满挑战，但也是学习TVM的关键步骤。

       学习TVM需要一定的前置知识，包括但不限于AI模型、深度学习框架、编程语言、Linux编程环境等。在学习过程中，遇到不懂的概念或难以理解的部分时，不要气馁，可以通过阅读文档、官方指南、参与社区讨论等方式寻求帮助。同时，实践是检验真理的唯一标准，通过不断的实践、调试和修改，可以更好地掌握TVM的使用与开发。

       总的来说，学习TVM既是一项挑战，也是提升AI领域技能的宝贵机会。通过合理规划学习路径，逐步积累知识和经验，每一个挑战都会成为进步的阶梯。在学习过程中，保持耐心和毅力，不断探索和实践，最终能够驾驭TVM，实现模型的高效部署与推理。

TVM源语-Compute篇

       本文探讨TVM源码中的计算相关(primitives)模块，深入讲解如何在非神经网络场景下，如基于张量的密集计算中，通过TVM的原生指令实现算法。通过分解计算与调度，TVM提供了一种灵活高效的并行计算框架。本文将首先通过向量相加（Vector Addition）实例，展示如何将算法数学表达式转化为TVM指令，实现输出矩阵的生成。接着，以矩阵乘法（GEMM）为例，说明TVM如何通过三层for循环来处理矩阵操作，并引入te.compute和te.reduce_axis等关键指令。进一步，同源码套通过简化卷积实现，解释了如何使用TVM DSL（数据描述语言）来处理多通道输入和输出特征图的卷积操作。最后，文章总结了TVM DSL的使用方式，强调其功能性编程风格，以及lambda表达式和reduce_axis在隐藏for循环细节、增强算法理解与优化后端性能方面的优势。

       在向量相加（Vector Addition）部分，我们定义数组长度n，两个数组A和B，通过lambda表达式将每个元素相加，存储到数组C中。TVM的te.compute指令用于指定输出结果的形状，lambda表达式则对应于循环逻辑，create_schedule构建出计算流程。利用tvm.lower将生成的schedule映射至IR（中间表示）上，展示与常规C代码相似的流程。

       矩阵乘法（GEMM）示例中，我们定义了矩阵A、B和C的维度，通过三层for循环实现矩阵乘法和加法。引入te.reduce_axis指令以优化循环结构，展示矩阵乘法运算的关键步骤和优化潜力。进一步，通过简化卷积实现，我们深入探讨了如何处理单通道输入图像和滤波器的卷积运算，解释了补零操作和使用te.compute处理多输入的实现方式。最终，总结了TVM DSL在表达计算逻辑、隐藏低级循环细节、优化算法性能方面的优势，以及其功能性编程风格对理解与优化算法带来的便利。

TinyML-TVM是如何驯服Tiny的（上）

       低成本、人工智能驱动的消费类设备的激增，激发了研究者对“裸智能”（低功耗、通常无操作系统）设备的兴趣。传统上，专家能在这些设备上运行某些模型，但为不同设备优化模型是个挑战，往往需要针对设备的手动优化。尤其是针对没有Linux支持的平台，没有可扩展的模型部署解决方案。因此，开发者通常需要创建一次性的定制软件堆栈，以管理资源和调度模型执行。

       尽管机器学习软件的优化并不是裸机领域特有的难题，它与GPU和FPGA等硬件后端的开发人员共同面对。TVM展现出了适应新硬件目标的能力，但在微控制器独特配置的挑战下，它之前还未能提供足够的支持。为解决这一问题，TVM扩展了微控制器后端，即µTVM（发音为“MicroTVM”），以在裸机设备上执行张量程序，并通过TVM内置的张量程序优化器AutoTVM自动优化这些程序。上图展示了µTVM+AutoTVM基础设施的概览。

       µTVM+AutoTVM的实际应用

       在讨论TVM/MicroTVM及其工作原理之前，我们先看一个实际应用示例。假设我们使用STMFZG板，它配备了一个强大的ARM Cortex-M7处理器，非常适合边缘人工智能应用。通过USB-JTAG端口将板连接至桌面，使用OpenOCD打开与设备的JTAG连接，从而通过µTVM使用设备无关的TCP套接字控制M7处理器。在桌面上运行TVM代码，执行CIFAR-分类器，如完整脚本所示：

       在这个设置中，性能表现可能不佳，但AutoTVM提供了一线生机。通过为设备编写调度模板并进行一轮自动调整，可以显著提升性能。具体来说，只需替换原始构建行：

       然后替换为：

       这样，结果将显著提升，性能大约提升2倍，与CMSIS-NN版本5.7.0（commit ab7c9a）相比，后者是一个手工优化的ML内核库。

       µTVM的性能表现与CMSIS-NN模型相比较具竞争力，但工作才刚刚开始，利用TVM的优化特性还有更多空间。通过加速密集/全连接运算（如密集/全连接操作）并利用TVM的模型特定量化和运算符融合功能，可以进一步优化性能。µTVM与TVM能够协同工作，发挥最佳性能。

       µTVM的设计理念

       µTVM旨在满足设备最低公分母的要求，只需提供设备的C交叉编译器工具链、读/写设备内存的方法、设备内存布局和体系结构特征的规范以及为设备准备函数执行的代码段。大多数裸机设备都支持C和JTAG，因此（1）和（2）通常是免费的。此外，（3）和（4）要求通常较小。例如，STMF系列板的（3）和（4）示例如下：

       µTVM基础设施和设备runtime的构建是为了满足这些需求，正努力通过支持常见的开源runtime平台（如mBED OS）来简化编译和链接过程。

       µTVM的设备会话

       考虑到微控制器的网络特性，引入了微会话的概念，它稍微偏离了标准的TVM代码。µTVM中的每一项功能都依赖于与目标设备的开放会话。在第一个代码片段中，一行代码偏离了规范，即这一行：

       通过这个with块内的每一行，都可以调用µTVM中的函数，上下文是device_config指定的设备。这条线背后做了很多工作，让其拆分如下：

       首先，它初始化与设备的连接，使用指定的任何通信方法（通常是OpenOCD）。然后使用指定的交叉编译器交叉编译µTVM设备runtime。最后，主机为编译后的二进制文件分配空间，并使用打开的连接将二进制文件加载到设备上。

       由于runtime现在位于设备上，自然需要一些函数来运行它。

       模块加载

       TVM的核心抽象之一是模块。模块为特定设备/ runtime目标存储一组相关函数。考虑到微控制器通常没有操作系统，µTVM需要额外的工作来维护这种高级抽象。跟踪创建和加载µTVM兼容模块的过程：

       假设有一个微型会议打开设备，并实现二维卷积的TVM调度。如果想把它加载到微控制器上，需要将C代码发出。为了实现这一点，只需要设定目标tvm.build或relay.build，例如：

       然后，通过µTVM基础设施中的核心功能运行它：create_micro_mod：

       这样，交叉编译模块中的C源代码，为生成的二进制文件分配空间，然后将二进制文件的每个部分发送到设备上分配的插槽中。一旦模块二进制文件在设备内存中处于合适的位置，二进制文件中的函数指针将被修补，使模块能够在设备runtime访问帮助函数（例如，分配草稿行）。

       加载内核后，可以获取卷积函数的远程句柄，如下所示：

       张量加载

       如果要调用运算符，首先需要一些张量作为参数：

       然后，根据其数据类型（例如int8、float等）和形状，计算每个张量的字节大小，主机在设备堆上分配内存区域。接着将张量的数据加载到分配的区域中。

       函数调用

       运算符执行可能是系统中最复杂的部分。为了简化表示，我们首先讨论严格执行（运算符一被调用就立即执行），然后是延迟执行（只有在需要运算符的结果时才执行运算符），这是系统的实际工作方式。

       严格执行

       调用函数时，输入和输出张量都作为参数传递，这就是目标传递样式：

       考虑到这些张量已经在设备上分配，只需要向设备发送元数据（设备地址、形状和数据类型），设备就能知道使用哪个驻留张量。下面显示了一个名为“runtime”的函数调用。在构造这个表示之前，需要将元数据序列化到设备上专门为此目的而存在的arguments部分中。

       µTVM会有一个全局UTVMTask实例，从主机端写入该实例。一旦写入任务，runtime就拥有了执行函数所需的一切，可以在runtime的入口点开始执行。runtime执行一些轻量级初始化，运行运算符，然后将控制权返回给主机。

linux本地clion调试TVM源码环境搭建

       首先，从网上下载TVM源码和LLVM，然后解压LLVM文件。

       接着，使用Clion打开TVM源码以CMake工程形式，确保在CMake选项中配置了解压后的LLVM路径。

       在成功加载CMake工程后，进行编译操作，点击工具栏上的编译按钮，编译结果会生成一个动态库文件，如libtvm.so。

       若遇到编译错误提示“unrecognized command line option ‘-fuse-ld=lld”，检查并升级gcc版本以解决此问题。

       仅需编译TVM代码即可开始调试工作，无需额外编译其他组件。

       准备Python代码执行环境，调整环境变量，确保PYTHONPATH指向TVM源码中的Python包路径，同时设置LD_LIBRARY_PATH指向动态库生成路径。

       尝试运行自编写的Python脚本，验证环境配置是否正确。

       为了调试C++源码，创建一个CMake应用，例如命名为cppEntrance，配置程序参数为待调试的Python脚本路径，并在环境变量中保持与Python脚本相同的设置。

       找到对应Python接口的C++代码入口，设置断点，启动cppEntrance调试，即可进入TVM的C++代码调试。

       对于查找TVM接口对应的C++代码入口，除全局搜索外，可能存在其他方法或工具。欢迎在评论区分享您的经验或建议。

编译TVM遇到 collect2: fatal error: cannot find 'ld'

       在尝试编译TVM源码时，我在Ubuntu.上遇到了一个错误，具体是"collect2: fatal error: cannot find 'ld'"。这让我感到困惑，因为我已经确认ld已经被安装。

       通过google搜索，我检查了ld的安装路径，发现它确实存在。我尝试了重新安装binutils，使用命令"sudo apt install --reinstall binutils"，但问题仍然没有解决。

       我重新思考问题，发现cmake编译时使用的链接器是lld，而不是官方推荐的ld。这可能是问题的关键。为了解决这个问题，我安装了lld（使用命令"apt-get install lld"），然后重新编译，这次成功了。

       TVM建议使用llvm进行编译，因此链接器使用的是lld。这次的经验让我意识到，确认链接器版本和正确配置cmake参数对于编译过程至关重要。通过这个过程，我学到了如何在遇到编译问题时，系统地排查并解决问题。

TVM 自底向上（二）：TIR 的概念和编译原理

       在深入探讨TVM中的编译过程与中间表示（IR）时，特别是TIR（Tensor IR）的概念及其编译原理，本节将重点聚焦于如何将神经网络模型转化为硬件源代码，以帮助读者更深入地理解这一复杂过程，并找到学习TVM的乐趣。

       TIR是TVM中最接近目标硬件的数据结构，无论前端模型（如pytorch、tensorflow或ONNX）经过了哪些转换，最终在被编译为目标后端代码前，都需要被转化为TIR。TVM的编译流程中，TIR起着核心作用，其位置如图所示。

       在TIR的实现中，抽象语法树（AST）扮演着关键角色。AST是一种通用的数据结构，用于表示任何编程语言的语法结构。它由节点组成，每个节点代表一种语言元素，如变量、函数调用或控制结构。在TIR中，AST为编译为不同硬件（如C++、CUDA、LLVM IR等）的代码提供了一个通用的结构。

       通过将AST转换为源代码（CodeGen过程），TIR能够解决神经网络推理计算中遇到的两个主要问题：首先，它能够表示深度学习算子（如卷积、池化、ReLU）和控制结构（如min、max、if-else），这些算子和控制结构都基于基本的数学运算。其次，TIR的通用性使得加速逻辑可以被抽象化并应用于各种硬件架构，从而实现跨平台的加速。

       TVM中的关键概念包括：IRModule、PrimFunc和CodeGen。IRModule是TVM中最小的编译单元，用于封装待编译的TIR和其他中间表示。PrimFunc封装了完整的AST，作为IRModule的API，对应生成.so库的函数入口。CodeGen负责将AST转换为目标硬件的源代码，本质上是一个树遍历迭代器。

       TVMScript提供了一种简化TIR AST开发的方法，它利用Python AST（Python的语法树实现，如Cython使用），允许直接使用Python语法编写AST，从而简化了TIR的开发流程。TVMScript还支持双向转换，即可以从TIR AST生成TVMScript，也可以从TVMScript解析回TIR AST。

       通过调用tvm.build函数，可以将IRModule编译为目标代码并运行，该过程根据所选的目标（如CPU、GPU或LLVM IR）选择适当的CodeGen。对于不同的目标，CodeGen过程涉及将TIR AST转换为目标硬件的源代码，然后使用相应的编译器生成可执行文件。例如，对于C++目标，CodeGen过程包括TIR到C++源代码的转换，而CUDA目标则涉及TIR到CUDA源代码的转换。

       最后，本节概述了使用TVMScript编写TIR AST和调用适当CodeGen编译为源代码的完整流程，并强调了其他相关章节的内容。通过了解这些概念和原理，读者能够更深入地理解TVM编译过程的内在机制，从而为探索和应用TVM提供坚实的基础。

在RK GPU 安装TVM

       在RK GPU上安装TVM

       RK搭载的ARM Mali G GPU，具备约GFLOPS的单精度运算能力，并支持OpenCL 2.1。本次将通过编译安装带有OpenCL功能的TVM，并在RK上运行模型。

       配置OpenCL编译选项是关键步骤。首先，在RK系统中使用`sudo find / -name *mali.so`查找libmali.so，这是包含OpenCL相关操作的动态库。然后，在build/config.cmake文件中，将USE_OPENCL项目填写为libmali.so的地址。

       接着，将OpenCL头文件目录CL放置于`/usr/include`中，并将路径添加到`~/.bashrc`的PATH变量中。执行命令于tvm的build目录下进行编译。

       在完成编译后，需先卸载已有的TVM版本。具体操作为：通过`pip3 uninstall tvm`命令进行卸载，接着进入源码目录`apache-tvm-src-v0..0/python`，删除指定的目录与文件。之后，进入tvm目录继续清理，确保环境整洁。

       安装TVM的过程如下：进入`apache-tvm-src-v0..0/python`目录执行相关安装命令。最后，通过验证步骤确认安装正确，查看输出结果，以确保TVM在RK GPU上运行正常。

       参考官方文档以获取更多详细信息：tvm.apache.org/docs/install/