1.LLVM(MLIR)安装编译
2.利用TPU-MLIR实现LLM INT8量化部署
3.MLIR Operation 与 Op
4.编译技术入门与实践之LLVM概述及环境构建
5.OpenAI/Triton MLIR 第零章: 源码编译
6.MLIR多层中间表示——用MLIR构建编译器(下)
LLVM(MLIR)安装编译
本文旨在为有兴趣自行安装和编译 LLVM(利用 MLIR 作为后端输出的源码主要方式)的读者提供一份详细指南。在实际操作过程中,解析可能会遇到一些理解上的源码偏差,欢迎指正。解析由于目标是源码能在 x 和 RISCV 上运行,所有配置均基于 i7-H 笔记本,解析小程序 商城 后台 源码运行 Ubuntu . LTS 系统。源码
以下是解析编译配置的步骤:
第一步:下载 LLVM 的源码。确保已安装 git,源码若未安装,解析请执行 sudo apt-get install git。源码创建名为 LLVM 的解析文件夹存放 LLVM 源码,并将源码文件夹命名为 llvm-project。源码接着,解析通过 git 下载 LLVM 源码。源码
第二步:建立用于 LLVM 编译的文件夹。为了区分编译产生的文件和源文件,建立名为 build 的文件夹。在教程中,每段代码都以 cd 到主文件夹,然后进入工程文件夹的方式进行,便于理解。
第三步:进入 build 文件夹,完成编译配置。此过程大致分为如何编译、编译什么、为谁编三个部分。具体参数如下:
如何编译:指定编译器类型、线程数及目标地址。例如,ocr助手源码使用 -DLLVM_PARALLEL_COMPILE_JOBS=### 设置并行编译工作数,使用 -DCMAKE_INSTALL_PREFIX=*** 指定安装路径,使用 -DLLVM_CCACHE_BUILD=### 选择是否使用 ccache。选择 C 和 C++ 编译器,如 -DCMAKE_C_COMPILER=### 和 -DCMAKE_CXX_COMPILER=###。启用 LLD 作为链接器以提高效率,可通过 -DLLVM_ENABLE_LLD=ON 实现。
编译什么:设置编译版本类型,如 Debug、Release 等,使用 -DCMAKE_BUILD_TYPE=###。同时,通过 -DLLVM_ENABLE_PROJECTS=### 配置需要编译的子项目。
为谁编:指定目标平台,如 x 和 RISCV,使用 -DLLVM_TARGETS_TO_BUILD=###。可选平台包括但不限于:AArch、AMDGPU、ARM、AVR、BPF、Hexagon 等。
注意:在完成编译配置后,执行编译命令。在遇到可能的问题时,检查错误信息并根据需要调整参数。最后,根据实际需求进行文件路径、编译选项等的hsf源码阅读调整。
以上步骤和参数配置将帮助您成功安装和编译 LLVM,满足在 x 和 RISCV 上运行的需求。通过本文提供的指南,希望能为您的项目开发提供便利。如有任何疑问或需要进一步的帮助,请随时提问。
利用TPU-MLIR实现LLM INT8量化部署
在年7月,我们已成功将静态设计应用于ChatGLM2-6B在BMX单芯片部署,采用F量化模式,模型大小为GB,平均速度为3 token/s。为提升效率与降低存储需求,我们进一步对模型执行了INT8量化部署。
传统TPU-MLIR的INT8量化方案并不适合LLM。这主要是由于LLM中PTQ校准或QAT训练成本过高,一轮校准可能需1-2天,且量化误差导致模型精度大量损失。基于此,我们沿用了ChatGLM2的W8A策略,对GLMBlock中Linear Layer权重进行per-channel INT8量化存储,运算时反量化至F,以确保精度损失几乎为零。
在编译器的Top至Tpu层lowering阶段,TPU-MLIR自动替换MatMul算子,将权重矩阵切分为W8AMatMul,以区分具有不同矩阵输入的算子。以ChatGLM2中某个MatMul算子为例,量化后权重从MB减至MB,额外的HPV溯源码Scale使用了0.MB存储,实现近一半的存储空间节省。相关源码可在TPU-MLIR仓库查询。
性能提升主要源于W8AMatMul后端算子优化。TPU架构下,W8A的计算过程分为5步,通过GDMA与BDC指令并行执行数据搬运与运算,将Local Memory分为两部分,确保效率。当左矩阵数据量较小时,性能瓶颈在于右矩阵数据加载,W8A量化减少数据搬运总量,额外运算时间被覆盖,性能影响可忽略。
从LLM角度看,推理流程包括prefill与decode。prefill阶段输入词向量补位至最大文本长度,decode阶段固定取前一轮生成的token作为输入。因此,prefill阶段GLMBlock接收数据量大时,W8A性能提升有限,而decode阶段$L_{ row}$恒为1,能实现显著性能提升。
应用W8A量化后,ChatGLM2-6B整体性能得到优化。具体结果展示如下:
MLIR Operation 与 Op
MLIR的核心组件之一是Operation,它是Dialect中的基本语义单元,就像方言中的词汇一样,构成了整个代码表示的基石。深入理解MLIR的源码输出 时钟运作,关键在于掌握Operation类和Op类的运作机制。
Operation类是MLIR中的核心数据结构,它在 mlir/include/mlir/IR/Operation.h 中被定义,为操作提供了丰富的接口,包括静态create函数,用于创建不同类型的Operation实例。例如,当我们创建一个FuncOp时,会先通过Operation::create调用,再通过cast转换为具体类。此外,Operation还负责操作数、结果、属性和blocks的管理,以及操作的遍历等。
Op类则是所有具体Operation的基类,它在mlir/include/mlir/IR/OpBase.td中定义,这个文件实际上是TableGen语言编写的。在定义新的Dialect时,开发者需要在Ops.td文件中引入OpBase.td,基于这个基类扩展出符合特定Dialect需求的Op类。
整个过程涉及到TableGen工具的编译和链接,将这些定义转化为编译时可用的代码。通过理解并操作这些基础类,开发者可以构建和优化复杂的编译逻辑。更多关于Op类的具体细节和功能,还需进一步研究mlir的源码和官方文档。
编译技术入门与实践之LLVM概述及环境构建
本系列旨在记录学习过程和知识总结,便于后续交流。我专注于智能芯片研究,编译器设计是实践中的常见挑战,最近实验中涉及LLVM pass,处理源代码到数据流的转换。
LLVM是一个集成了模块化技术的编译器项目,非传统虚拟机概念,全称为LLVM项目。最初由伊利诺大学发起,目标是创建基于SSA策略的现代化编译策略,支持多种语言的静态和动态编译。它包含众多子项目,如LLVM core、Clang、LLDB等,广泛应用于学术和商业领域,以其通用性、灵活性和可重用性为特点。
LLVM的核心子项目包括:LLVM core提供源码和目标代码隔离的优化器,支持多种CPU;Clang是C/C++编译器,包含自动code检查工具;LLDB是高效调试器,基于LLVM核心和Clang;libc++与libc++ ABI提供C++标准库;compiler-rt提供底层代码生成优化;MLIR构建可扩展的编译器架构;OpenMP支持OpenMP在Clang中的使用;Polly进行本地化优化和并行化;libclc开发OpenCL标准库;klee则提供符号化虚拟机以检测bug。
实验准备部分,LLVM项目包含一系列工具和库,如汇编器、反汇编器等,用于处理LLVM中间表示和目标文件的转换。获取源代码时,需要明确硬件平台和软件环境。在Ubuntu系统中,可能需要升级cmake、安装openssl等。构建LLVM和Clang时,遵循官网指南,配置后执行make或ninja命令进行编译,并可能需要解决一些编译错误。
OpenAI/Triton MLIR 第零章: 源码编译
本文旨在深入探讨开源AI项目OpenAI Triton MLIR,着重介绍Triton作为编程语言与编译器在GPU加速计算领域的应用与优化。Triton为用户提供了一种全新的方式,通过将其后端接入LLVM IR,利用NVPTX生成GPU代码,进而提升计算效率。相较于传统CUDA编程,Triton无需依赖NVIDIA的nvcc编译器,直接生成可运行的机器代码,体现出其在深度学习与数据科学领域的高性能计算潜力。Triton不仅支持NVIDIA GPU,还计划扩展至AMD与Intel GPU,其设计基于MLIR框架,通过Dialect支持多样化后端。本文将从源码编译角度出发,逐步解析Triton的设计理念与优化策略,为研究编译技术和系统优化的工程师提供宝贵资源。
首先,需要访问Triton的官方网站,克隆其官方代码库,以便后续操作。构建过程涉及两个重要依赖:LLVM与pybind。LLVM作为Triton的核心后端,通过将高级Python代码逐步转换至LLVM IR,最终生成GPU可运行代码,体现了其在计算优化领域的优势。pybind组件则用于封装C++/CUDA或汇编代码,实现Python DSL与高性能组件的无缝集成。
接下来,将LLVM与pybind分别编译安装,通过手动配置指定路径,确保编译过程顺利进行。LLVM的安装对于基于Triton进行二次开发的工程师和研究人员至关重要,因为它为Triton提供了强大的计算基础。在特定的commit ID下编译Triton,确保与后续版本兼容。
在编译过程中,配置pybind同样至关重要,它允许用户通过Python API调用高性能组件,实现自动化生成高性能算子。完成编译后,生成的.so文件(libtriton.so)为后续Triton的Python接口提供了支持。
将libtriton.so移动至triton/python/triton/_C目录下,确保Python路径正确配置,实现无缝导入与调用。通过简单的import triton命令,即可开启Triton的开发之旅。验证Triton性能,可以选择tutorials目录下的示例代码,如-matrix-multiplication.py,通过运行该脚本,观察Triton在GPU上的性能表现。
Triton在NVGPU上的成熟映射路线,从抽象的Python DSL到贴近GPU层面的IR,最终生成高效机器代码,体现了其在高性能计算领域的优越性。Triton未来的发展蓝图将支持更多前端语言,对接不同硬件厂商的硬件,实现高效映射,满足多样化计算需求。
MLIR多层中间表示——用MLIR构建编译器(下)
在构建编译器的过程中,将方言翻译到LLVM的策略为走向代码生成的关键步骤。此过程涉及到将源代码的表示转换为执行代码的中间表示,以便进一步优化并最终生成可执行文件。
在MLIR中,实现这一目标的方法是利用LLVM方言,这是一种预定义的模式集合,允许从源语言到目标语言的转换。通过连接现有的转换框架和组件,可以建立一个完整的端到端系统,用于从源语言编译至可执行代码。
系统的核心构建包括定义转换目标、运算转换和类型转换。转换目标说明了哪些运算是合法的以及在何种情况下,运算转换负责将非法运算转化为合法形式,而类型转换则规范了非法类型如何被合法化。在方言转换中,有两种模式可供选择:部分模式允许不是所有输入运算都必须对目标合法化,而完整模式则要求所有输入运算都必须对目标合法化。
通过定义转换目标和收集运算转换模式,构建了一个可以将Toy语言转换为可执行代码的系统。在过程中,需要处理那些现有运算可能无法合法化的特殊情况,允许在不了解整个IR的情况下转换已知非法运算的子集。
当转换系统成功地应用在函数上时,它会尝试将非法运算转换为合法运算,如果任何非法运算未能成功转换,则转换过程会失败。在实际应用中,示例代码展示了如何从Toy语言转换到Affine语言,并执行了简单的计算操作,最终输出结果并返回。
在完成方言转换后,可以将MLIR LLVM方言导出为LLVM IR,实现从LLVM方言到LLVM IR的映射。此过程涉及将MLIR模块转换为LLVM IR,以便进一步在LLVM环境中执行和优化。
整个过程展示了从源语言到可执行代码的高效转换路径,利用MLIR和LLVM的集成能力,构建出强大的编译器系统。此系统不仅能够实现语言之间的转换,还能在转换过程中进行代码优化,最终生成高性能的执行代码。