LLVMCon-US-2025-笔记-全-

LLVMCon US 2025 笔记（全）

001：LLVM 预合并测试新系统 🚀

在本节课中，我们将学习 LLVM 项目全新的预合并测试系统。这个系统旨在提升代码提交前的测试速度、可靠性和易用性，帮助开发者更高效地发现和修复问题。

系统概述与背景

上一节我们介绍了课程主题，本节中我们来看看新系统的背景和目标。

预合并测试在开发者创建或更新拉取请求时，会自动触发一系列检查和测试。这些测试包括在不同平台上运行 Ninja 检查测试。Linux 和 Windows 测试由 Google 在 Google Cloud Platform 上运行和支持。最近，亚马逊 AWS 上也启动了针对 AArch64 架构的测试，这些测试也能从我们的工作中受益。

预合并测试对开发者非常有用，因为它帮助你在代码实际提交到上游之前发现和识别错误。特别有用的是，它在一系列平台上运行测试，而你个人可能无法访问所有这些平台。当预合并测试通过时，你对你的更改不会意外破坏上游内容更有信心。

LLVM 的预合并测试自 2019 年就已存在。旧系统运行在 Buildkite 上，速度慢、不稳定、不可靠，结果不一定正确，并且维护不佳。我们认为这不是一个理想的状态，因此我们一直在开发一个新系统。今年五月，我们正式推出了新系统。它运行在更快、更专用的机器上，并且由于具备自动扩缩容功能，效率更高。我们使用的机器数量会根据实际运行的测试数量动态增减。

我们还在工作时间设有专门的待命支持轮换，因此有人实际监控预合并测试系统。如果出现问题，我们会立即识别并处理，确保系统得到良好支持和修复。

我们还为系统添加了指标和性能跟踪功能，并在仪表板上展示。现在我们可以看到使用了多少台机器、运行了多少作业、作业耗时、任务在真正启动前排队了多久等信息。当然，我们也将这些信息用于待命轮换的跟踪。

我们于七月初将 libc++ 的预合并测试迁移到了新系统，这解决了 libc++ 社区在测试中遇到的许多问题，libc++ 的维护者对我们的系统特别满意。

这个新预合并测试基础设施的总体目标是：速度、可靠性和易用性。

我们理想的目标是让 80% 到 90% 的测试在 30 分钟或更短时间内完成。我们希望确保没有误报的错误报告，即如果测试系统告诉你你的拉取请求破坏了某个特定测试，那么你的拉取请求确实破坏了该测试，而不是因为测试不稳定或主线上的其他故障。此外，我们还希望让你能非常容易地找到失败原因并理解其根源，以便更快地调试和修复。

新功能：失败摘要与自动合并

上一节我们了解了系统的目标，本节中我们来看看两个提升体验的新功能。

以下是两个我们添加的、我们认为有助于实现上述目标的新功能。

第一个是失败摘要。这里有一个显示某些测试失败的拉取请求的屏幕截图。你可以点击三个点，这会弹出“查看详情”选项。点击“查看详情”后，会跳转到失败详情页面。在“构建和测试”部分，你可以看到所有来自脚本日志和 Ninja 日志的日志行，你可以滚动浏览这些内容来查找所有构建失败。但这些日志可能长达数千行，滚动浏览非常耗时，并且很难找到你真正需要的信息。

因此，我们添加了“摘要”按钮。点击“摘要”后，会跳转到针对不同架构的测试摘要页面。向下滚动到你感兴趣的架构，例如 Linux，我可以立即看到哪些测试在 Linux 上失败了。我甚至可以点击其中一个测试来获取实际的测试错误信息。这比试图滚动浏览数千行 Ninja 构建日志要快得多，也容易得多。

我想指出的另一个功能是，我们增加了无需等待和监控测试的能力。我们提供了两种方式，让你不必坐着等待测试完成。如果你确信想要合并你的更改，并且不需要等待测试结果（例如你正在进行一个非常快速和简单的还原操作），你可以点击“无需等待要求即可合并”，然后立即合并你的更改。

另一方面，如果你确实关心测试结果，但仍然不想坐着等待拉取请求的测试完成，你可以点击“启用自动合并”。这将允许测试运行完成，如果测试通过，它会自动将你的更改推送到代码库，你无需自己操作。但如果测试失败，则不会进行合并。这些是我们添加的一些功能。

性能优化与改进

上一节我们介绍了提升用户体验的功能，本节中我们来看看系统在性能方面的具体改进。

现在我们在性能方面的一大改进是，我们建立了一个 Grafana 仪表板来跟踪所有这些指标，这样我们就能立即看到是否有性能回退。待命轮换的人员可以查看仪表板，了解发生了什么，例如是否有任务长时间排队。我们为 libc++ 仪表板也设置了同样的支持。此外，我们还支持查看随时间变化的趋势，以判断我们是否做出了真正有意义的改进。这些图表中有很多噪音，因为如果有人修改了一个 ADT 头文件，缓存命中率就会突然变化，或者当天人们提交的代码内容也会带来很多噪音。

尽管如此，我们仍然可以看到一些显著的性能改进。例如，在七月中下旬，我们显著改进了编译缓存的架构，这是一个相当大的改进。当我们迁移时，Buildkite 系统会在本地缓存所有内容，因为它会在多次运行之间保留一些状态，但这并不是最佳方式。根据 GitHub 的工作方式，我们必须使用一个干净的环境。我们曾使用 GitHub 缓存，但效果不佳，因为它无法在拉取请求之间有效地共享编译缓存。因此，我们将其设置为使用 Google Cloud Storage，将所有内容存储在一个云存储桶中。这最终显著提高了缓存命中率。

现在，大多数拉取请求（假设你没有修改某些公共头文件）的缓存命中率约为 99% 以上，这导致整体测试时间减少了 30% 到 35%。现在编译时间通常只占很小一部分。因此，测试时间现在在预合并测试的实际运行时间中占主导地位。我们一直在努力将测试迁移到 LLVM 的内部 shell，这减少了一些进程启动开销。这在 Linux 端带来了 10% 到 15% 的性能提升。Windows 默认已经在使用它，因为那里的进程启动开销更为显著。所有使用它的主要测试套件（除了 compiler-rt 之外的所有测试）目前都默认使用它。我有一个补丁堆栈需要提交以使 compiler-rt 正常工作，clang 方面还有一些问题我们正在解决。这样，你在预合并测试中就能获得性能提升，在本地运行 check-llvm 时也能获得，这对本地开发周期甚至更有益。

我们还致力于在 Windows 上改用 clang。如果你获得了高缓存命中率，那么具体使用的编译器速度并不太重要。但当你需要重新构建大量内容时，编译时间就很重要了，这里的“长尾”效应可能相当显著。在 Windows 上使用 clang 而不是 MSVC 导致编译时间减少了 40% 到 50%。我认为这主要是因为工具链的构建方式，实际上 Windows 上的发布版工具链是专门为构建 LLVM 进行了 PGO 优化的，这肯定有帮助。

减少测试不稳定性与未来展望

上一节我们讨论了性能提升，本节中我们来看看如何减少测试的不稳定性。

正如之前提到的，我们的目标是确保在你的拉取请求上报告的失败确实是由你的拉取请求引起的。因为 LLVM 是一个大项目，存在不稳定的测试和由于“空中碰撞”导致的失败（即使预合并测试通过，但当你提交时，它与之前提交的更改发生冲突）。

我们设置了 Buildbot，目前它只连接到暂存的构建主服务器，这使我们能够轻松查看主线代码的当前状态。然后，我们将这些数据以及来自拉取请求的失败信息输入到我们称为 “预合并顾问” 的系统中。这个系统目前大部分已实现，我们只需要做一些收尾工作，然后实际在摘要视图中（可能也在拉取请求的评论中）展示这些信息。

如果你的拉取请求遇到了已知的不稳定测试或主线已经失败的测试，系统会将其标记为“主线失败”，而不是“你的拉取请求内失败”。这样，预合并测试就会通过，并且我们消除了所有误报的测试报告。同时，它可能会给你一些链接，指向实际失败的后置提交运行，如果你有兴趣，可以确认这确实是一个不稳定的测试。

不稳定的测试通常很难发现，除非你从整体上查看所有基础设施。既然我们拥有这些信息，我们希望找出如何轻松地向社区公开这些信息的方法。以前我遇到过不稳定的测试，通常是有人在 Discord 上私信告诉我。拥有一个更好的方式来处理这个问题会很有帮助。最终，理想情况下是消除所有测试的不稳定性。

本节课中我们一起学习了 LLVM 全新的预合并测试系统。我们了解了其背景、目标、提升用户体验的新功能（如失败摘要和自动合并），以及通过改进缓存架构、迁移到内部 shell、使用 clang 等带来的显著性能提升。最后，我们还探讨了旨在减少误报和测试不稳定性的“预合并顾问”系统。这些改进共同为 LLVM 开发者提供了更快、更可靠、更易用的代码提交前测试体验。

002：2025年更新

在本节课中，我们将回顾Clang编译器及其社区自2024年10月以来的最新进展。我们将涵盖C和C++语言标准支持、缺陷报告修复、社区贡献模式以及用户反馈等多个方面。

C语言标准支持进展

上一节我们介绍了课程概述，本节中我们来看看Clang在C语言标准支持方面的具体工作。过去一年，团队主要致力于C23和C2Y（C语言的下一个标准草案）功能的实现。

以下是Clang各版本在C语言方面的主要新特性：

  • Clang 20：改进了枚举支持，现在可以像C++一样为枚举指定固定的底层类型（如 enum E : long long），这是一个已被广泛支持了数十年的扩展。
  • Clang 21：
    • 实现了C2Y的 countof 特性，用于确定数组的大小。
    • 添加了八进制字面量的新前缀 0o（类似于已有的 0x 和 0b）。
    • 改进了标签兼容性。C语言没有C++那样的“单一定义规则”，而是基于类型兼容性。新规则允许在同一翻译单元内重新定义结构体，只要所有成员和内存对齐方式完全相同。虽然这可能导致一些长期存在的bug，但它能更好地支持匿名结构体的使用。
  • Clang 22（进行中）：
    • 实现了命名循环（named loops），允许通过标签从多层嵌套循环中跳出或继续，使控制流更清晰。
    • 正在实现 defer 技术规范。该特性类似于C++的RAII，但适用于没有构造/析构函数的C语言。它允许将一段代码的执行推迟到当前作用域结束时。

此外，在所有Clang版本中，团队都在根据C标准委员会的“Earthly Demons”提案，将一些过时且无用的未定义行为（UB）转换为约束违规，从而提供明确的诊断信息。

C语言标准支持度概览

以下是截至2025年的C语言各标准支持完成度（文本为2024年数据，用于对比进展）：

  • C99：100%完成（过去一年无变化）。
  • C11：82%完成（未知部分已全部解决）。
  • C23：79%完成（较2024年的76%有所提升）。
  • C2Y：约81%完成（注意：C2Y尚未定稿，其内容在不断变化）。

C++语言标准支持进展

上一节我们了解了C语言的进展，本节中我们来看看C++方面的更新。Clang在C++标准支持上也投入了大量工作。

以下是Clang各版本在C++语言方面的主要新特性：

  • Clang 20：
    • 实现了C++23中基于范围的for循环的生命周期扩展，消除了一个常见的由引用导致的未定义行为，减少了程序出错的可能。
    • 改进了常量表达式支持，包括 constexpr 中的未知指针/引用以及 constexpr 替换 new。
    • 开始支持C++26的 [[unsequenced]] 和 [[reproducible]] 属性。
  • Clang 21：
    • 继续推进C++26支持，包括在if语句等条件中使用的结构化绑定。
    • 为libc++实现了平凡可重定位性（trivial relocatability）特性。请注意，此特性可能因即将到来的C++会议讨论而调整。
  • Clang 22（进行中）：
    • 当前阶段主要以修复bug为主，因为C++是一门极其复杂的语言。
    • 预计未来将重点投入契约（Contracts）和反射（Reflection）两大特性。这两个都是多年期的大型项目，几乎不可能在Clang 22中完成。目前已有实验性的分支，相关贡献者正计划逐步将其合并到上游。

C++语言标准支持度概览

以下是截至2025年的C++语言各标准支持完成度（文本为2024年数据）：

  • C++20：95%完成（较2024年的92%提升，主要归功于模块工作的推进，模块是完成C++20支持的最后一块拼图）。
  • C++23：91%完成（较2024年的89%略有提升）。
  • C++26：约58%完成（较2024年的83%下降，这是因为C++26近期才功能冻结，新增了许多大型特性，属于正常波动）。

缺陷报告（Defect Reports）处理情况

上一节我们介绍了新特性的实现，本节中我们来看看对现有标准中错误的修复情况。缺陷报告（DR）是标准文档中的bug，实现它们对保持编译器一致性至关重要。

以下是缺陷报告的处理状态：

  • C语言缺陷报告：50%已实现，9%状态未知，2%明确不实现。
  • C++语言缺陷报告：21%已实现，54%状态未知，2%明确不实现。

“未知”状态占比很高，这意味着社区急需帮助来编写测试用例，以确定Clang是否已经支持了这些缺陷报告。标准委员会在制定新特性时，默认所有实现都已修复了所有缺陷报告。

从趋势来看，Clang 18-20在修复缺陷报告方面做得很好，但Clang 21有所下降，Clang 22至今尚未修复任何缺陷报告。趋势线开始转向负面，未来可能新增的缺陷会比修复的更多。

这项工作虽然不 glamorous，有时还很令人沮丧，但极其重要，社区非常需要帮助。

社区贡献分析

上一节我们讨论了技术工作，本节中我们来看看Clang社区的贡献模式。了解社区如何运作有助于我们更好地参与和优化流程。

工作时间分布

数据分析显示，Clang社区主要是“朝九晚五”型，大部分工作集中在周一至周五。周末的工作量虽然较少，但呈上升趋势。有趣的是，对比2023年和2024年的年底（圣诞节期间），2023年大家只在圣诞周休息，而2024年似乎整个12月工作量都大幅减少。

代码审查与PR合并速度

以下是关于Pull Request处理速度的数据分析：

  • 总体趋势向好：超过一周才关闭的PR比例在下降，这意味着社区在高效合并PR方面有进步。
  • 一个值得关注的现象：在一天内就关闭的PR比例迅速上升。原因可能是多方面的：供应链安全补丁增多、代码审查实践变化（如更多自检通过），或者审查质量下降。目前尚无定论，但值得关注。
  • 大多数PR在一天到一周内被处理，这是一个比较健康的速度。

贡献者与未合并的PR

贡献者几乎全部来自GitHub上标记为LLVM成员或贡献者的人。好消息是，首次贡献者的数量在增加，这是社区未来的希望。

然而，一个令人担忧的问题是：首次贡献者提交的PR几乎总是处于长期开放状态，无人审查。这是因为首次贡献者无法在UI上指定审阅者，系统也不会自动添加。这导致他们的PR很容易被忽视。社区需要更主动地关注这些PR，直到GitHub授权策略改善。

总体来看，处于开放状态的PR数量随时间在增加。这表明代码审查的带宽是Clang发展的瓶颈。我们需要更多专注的审阅者。

如果你有机会，请参与代码审查。你不需要是专家，即使只是指出代码风格问题，也是巨大的帮助。

用户生态与未来重点

上一节我们分析了社区内部，本节中我们来看看外部用户的情况和未来的工作重点。

用户使用的语言模式

根据Sourcegraph对开源项目的分析（数据不包含私有项目，且可能不完整）：

  • C++代码：90%使用C++17或更早的标准。
  • C代码：95%使用C11或更早的标准。
  • 最新标准：C++23和C23的占比分别仅为2%和1%，C2Y和C++26甚至未在图表中显示。

这强烈表明，确保现有代码的稳定运行对用户至关重要。 但这并不意味着我们应该忽略新标准，因为用户期望我们支持所有标准。这提醒我们，不必为了急于推出C++26等新特性而牺牲旧语言模式的代码质量。

用户反馈与社区挑战

从网络讨论和会议交流中，我们观察到：

  • 新特性受关注：defer技术规范获得了超过300个GitHub点赞，社区对此非常兴奋。
  • 安全特性受欢迎：边界安全检查、生命周期安全分析、各种消毒剂（Sanitizers）等安全改进工作持续受到用户高度赞赏。
  • 新标准尝鲜者：有用户在尝试C23和C++23并反馈问题，这对我们很有帮助，但绝大多数用户仍停留在C99和C++11。

社区内部面临的挑战和需要帮助的领域包括：

  1. 完成C++20模块：这是宣告C++20支持完成的最后障碍。模块维护者已公开求助，但响应不足。如果你对语言支持感兴趣，帮助完善模块将是对项目的巨大贡献。
  2. 应对大型新特性：契约（Contracts）和反射（Reflection）等特性需要巨大的投入，包括大量测试用例和文档。
  3. 编译时间与文档：用户普遍抱怨编译时间越来越长，以及文档分散且难以查找。我们需要偿还这些技术债，降低编译时间以保持竞争力，并大力改进文档体系。对于新手来说，参与文档改进是一个绝佳的入门方式，门槛较低但对用户影响巨大。

总结与问答摘要

本节课中我们一起回顾了Clang在2024-2025年间的进展。我们涵盖了C/C++新标准支持情况、缺陷报告修复的紧迫性、社区贡献模式的洞察，以及用户生态的现状。核心挑战在于代码审查带宽不足、大型新特性带来的可持续性压力，以及改善编译时间和文档的迫切需求。

以下是演讲后问答环节的摘要：

  • defer与cleanup属性：defer本质上是cleanup属性的语法糖，两者使用相同的底层机制，会自然地交互。
  • 支持未来C/C++特性的可持续性：演讲者对此表示担忧。像模块、契约、反射这样的大型特性实现和维护成本极高。虽然委员会可能认为这是在给编译器团队“找活干”，但一旦标准加入新特性，用户需求就会随之而来。最好的影响途径是参与C/C++标准委员会，从设计阶段就提出实现复杂度的考量。
  • PR处理速度提升的原因：数据显示PR合并速度确实改善了，可能原因是社区规模增长，提交和审核的PR都变多了，但缺乏更细粒度的数据来确定具体原因。

本节课中我们一起学习了Clang编译器在过去一年的发展全貌，从技术实现到社区生态，并明确了未来需要共同努力的方向。

003：Mojo的经验教训

在本教程中，我们将学习如何使用MLIR构建现代编程语言的前端，并以Mojo语言作为案例。我们将重点探讨Mojo中元编程的实现方式，以及如何利用MLIR的独特能力来支持复杂的参数化类型系统和编译时计算。

😊 概述：Mojo与元编程

Mojo是一种为异构、分布式和加速计算设计的通用编程语言。它建立在MLIR之上，旨在提供零成本抽象和安全性。Mojo的一个核心特性是其强大的元编程能力，这允许开发者将许多传统上由编译器处理的逻辑（如针对特定形状的代码特化、自动调优等）转移到库中实现。

上一节我们介绍了Mojo的基本目标和背景，本节中我们将深入了解其元编程的工作原理。

😊 Mojo元编程基础

Mojo的函数风格与Python类似，接受运行时参数。但Mojo还引入了编译时参数，我们称之为“参数”，它们用于参数化代码。

核心概念：

  • 运行时参数：在函数调用时确定值的普通参数。
  • 编译时参数：在编译时已知的参数，用 参数 关键字声明，为代码提供了“超能力”，例如保证完全展开的循环。

以下是一个简单的幂函数示例，展示了参数的使用：

fn power[base: Int, exp: Int]() -> Int:
    var result = 1
    for i in range(exp):
        result *= base
    return result

在Mojo中，不仅函数可以被参数化，任何结构体都可以。参数可以是类型、函数闭包甚至任意表达式。

😊 在MLIR中实现参数化IR

Mojo利用MLIR的属性系统来构建其参数化中间表示。理解这一点需要区分两个概念：

  1. 基础IR：承载运行时应用逻辑的基本IR构建块。
  2. 元IR：操作嵌套在其中的任何基础IR。

我们可以将定义参数视为在结构中“挖洞”，而实例化则是用正确类型的值“填充这些洞”。

定义参数化类型与操作

在传统方言中，定义类型（如静态大小的张量）可能会直接指定具体属性。而在参数化方言中，我们希望抽象掉具体细节，只约束其类型。

例如，一个参数化的张量类型可能将其大小定义为一个类型化的“洞”，仅在C++验证中确保其类型为索引类型，而不限制具体值。

操作也是如此，任何希望被参数化的操作属性或性质都可以进行相同处理。这是实现程序化代码生成的关键部分。

参数化函数与内联实例化

有了基础的IR构建块后，我们可以将它们封装到函数中。参数化声明（如参数化函数）将一段参数化IR封装在一组输入参数之上。

在函数体内，可以引用这些输入参数，并保证它们具有声明的类型。这允许我们使用不同的具体输入参数来复用该函数体。

此外，Mojo提供了典型的元编程指令，如 参数 if 和 参数 for，用于指定何时以及多少次内联实例化嵌套的参数化IR。

综合来看，参数化IR由两层组成：底层是带有许多“洞”的参数化操作，上层是参数层。这些参数被填入底层的洞中。MLIR的伟大之处在于，它允许元IR操作和基础IR操作自然地交织在一起，从而保持了方言的可组合性。

😊 参数计算与规范化

仅仅拥有参数引用还不够，我们还需要能够对这些属性进行计算。这带来了类型相等性判断的挑战。

在Mojo中，类型本身也是编译时值，因此存在依赖类型。判断两个依赖类型是否相等，实际上需要判断它们的参数表达式是否相等。

我们通过规范化来解决这个问题。将参数表达式规范化为一种标准形式，使得本应相等的表达式最终具有相同的标准形式。这意味着我们需要一个非常强大的表达式规范化器。

在实践中，我们将参数表示为类型化属性。属性在MLIR上下文中是唯一的，这使得参数相等性检查简化为指针相等性检查，无论表达式多复杂，都是常数时间。同时，它还减少了内存占用，并提供了清晰的打印表示。

即使参数表达式存在于上层的参数域中，它仍然可以引用底层用户声明的函数。这形成了一个完整的循环：不仅是元代码，代码也可以是元代码。用户可以在参数表达式中调用函数。

😊 实践优势与总结

我们进行所有这些工作的主要目的是实现真正强大的库。开发者可以构建高级抽象（如基于瓦片的编程模型）作为库，这大大简化了编程模型。

此外，调试器可以正常工作。如果你想调试元程序，只需在运行时调用它并单步执行。这与C++模板等机制有巨大区别。

参数化IR的另一个关键优势是，基于图的编译器和其他上层工具可以处理这种参数化的、可移植的表示形式，使得整个系统能够很好地组合在一起。

本节课总结：
我们一起学习了Mojo如何利用MLIR构建其现代语言前端，特别是其元编程系统。关键点包括：区分编译时参数与运行时参数；在MLIR中使用属性系统表示参数化IR，形成基础IR与元IR的双层结构；通过表达式规范化解决依赖类型的相等性判断；以及这种设计带来的强大库支持、简化编程模型和良好调试体验等实践优势。Mojo的经验表明，MLIR是构建支持复杂元编程和异构计算的现代语言的强大基础。

004：一个孵化器项目之旅 🚀

在本教程中，我们将学习ClangIR项目如何从一个独立的开源项目，经过社区孵化，最终成功上游化到LLVM主仓库的完整历程。我们将探讨孵化器的运作模式、社区协作的挑战与收获，以及上游化过程中的关键策略和经验。

概述 📋

ClangIR项目源于优化和分析C++代码的愿景。在传统的Clang编译流程中，AST（抽象语法树）层级过高，而LLVM IR层级又过低，缺乏对C++语义的丰富表达。ClangIR旨在填补这一空白，提供一个更丰富的中间表示（IR），以便进行更深入的语义分析和优化。

从RFC到孵化器 🏗️

上一节我们介绍了ClangIR项目的起源。本节中，我们来看看它是如何从一份RFC（征求意见稿）进入LLVM孵化器的。

项目时间线始于2021年底，最初在Mantta的GitHub上作为开源项目启动。经过约8个月的初步开发，团队向LLVM社区提交了RFC，寻求反馈。社区讨论的成果是成立了一个“C++ MIR前端工作组”，该工作组每月举行会议，持续了三年，汇集了行业专家共同探讨。

基于RFC的积极讨论，项目被接纳进入LLVM的“孵化器”。孵化器是LLVM GitHub组织下的一个独立目录，本质上是一个项目分支，为实验性项目提供了发展空间。

以下是孵化器的主要优势与挑战：

优势：

  • 可见性提升：作为LLVM伞形项目的一部分，能吸引更多关注和贡献者。
  • 基础设施支持：可以继承LLVM主项目的CI（持续集成）等基础设施，便于测试和代码规范检查。
  • 实验自由：由于审查者相对较少，在遵循LLVM实践的前提下，有更多空间进行实验和快速迭代。

挑战：

  • 非主线状态：对于下游公司或用户来说，依赖一个孵化器分支仍存在风险，且不如主仓库有吸引力。
  • 流程不明确：孵化器项目何时、如何能“毕业”进入主仓库，缺乏正式化的标准流程。
  • 同步负担：需要不断变基以跟上LLVM主项目的快速变化，维护成本高。

孵化器内的协作与成长 🤝

在孵化器中，项目的社区生态逐渐形成。贡献模式呈现出有趣的特点：初期有贡献爆发，随后趋于平稳，在2024年初因上游化进程而再次获得持续贡献。

贡献者来来往往是开源项目的常态。为了帮助新贡献者快速入门，项目将一些工作（如内置函数或内部指令的实现）设计为入门任务。这些任务通常只需10-15行代码和一个测试用例，让贡献者能快速获得成就感并理解代码库。

这种策略虽然有效，但也带来挑战：难以积累长期、深入的审查专家。项目贡献者“寿命”图显示，虽然有长期贡献的核心成员，但更多是短期贡献者。

一个积极的转变是，贡献来源从最初由Mantta主导，逐渐扩展到更广泛的社区，包括NVIDIA等公司的工程师、Google Summer of Code学生以及独立开源爱好者。这证明了社区协作的价值。

回顾孵化器阶段，尽管存在各种挑战和激烈讨论，但最终通过沟通找到了共同点，建立了合作关系，整体是一次非常积极的经历。

上游化：从零开始的艺术 🎨

随着项目在孵化器中逐渐成熟，团队于2024年初提交了上游化RFC，并获得了多家公司和项目的支持。2024年4月，第一个PR成功合并到LLVM主仓库。

上游化过程的一个关键策略是：由与原始孵化器开发不同的贡献者群体来主导上游化工作。这看似意外，实则带来了巨大好处。它类似于“重写《堂吉诃德》”——新贡献者必须深入理解每一行代码，仿佛第一次开发一样，从而带来了全新的视角和更深的理解。结果是，对代码有深刻理解的人数翻了一番。

上游化必须遵循LLVM的高质量标准：

  • 每个提交都必须功能完整、可测试，且足够小以便审查。
  • 贡献者需要彻底理解代码，能在代码审查中为其设计辩护。
  • 必须符合编码规范，保持提交历史清晰。

因此，团队将大约10万行代码，分解成小块，逐一上游。自2025年2月以来，项目保持了约每周1500行代码的稳定上游进度。

上游化策略与当前进展 📈

目前，上游化工作主要遵循经典的代码生成路径，专注于Linux/x86平台，以确保基本功能的正确性和可测试性。长远来看，ClangIR的潜力在于其“MLIR消费者”环节，它可以为异构计算、特定领域优化等打开新路径。

以下是推动进展的具体方法：

  • 测试驱动：使用LLVM测试套件中的“单源”测试作为工作清单，逐个解决阻碍编译的语言特性问题。
  • 对齐经典CodeGen：ClangIR的代码生成在结构上刻意模仿现有的LLVM IR生成路径（称为“经典CodeGen”），以降低学习成本。
  • 对比测试：在测试中同时输出ClangIR路径和经典CodeGen路径产生的LLVM IR，进行可视化对比，确保语义等效，并及时发现经典CodeGen的变更。

一个核心挑战是代码重复。ClangIR生成和LLVM IR生成的代码结构高度相似，仅因操作的值类型不同（CIR Value vs LLVM Value）而无法直接共享。团队正在探索使用模板或概念来抽象IR底层，以实现代码复用，这是长期可持续发展的关键。

总结与展望 🌟

本节课中我们一起学习了ClangIR项目从孵化到上游的完整旅程。

总结要点：

  1. 孵化器价值：为项目提供了可见性、协作平台和实验自由，是培育社区和验证想法的宝贵阶段。
  2. 社区力量：广泛的社区贡献是项目成功的基石，来自不同背景的贡献者带来了活力和深度。
  3. 上游化艺术：高质量的上游化需要将大代码库分解，遵循严格标准，并由新鲜视角进行“重写式”迁移，这反而加深了社区对代码的理解。
  4. 当前状态：ClangIR已具备一定语言支持，适合开发者开始在其上构建优化实验。项目正稳步推进，致力于实现与经典CodeGen的语义对齐和未来代码复用。

ClangIR的上游化不仅是代码的迁移，更是一个社区构建、知识传递和工程卓越的过程。它为未来更强大的C++编译优化能力奠定了基础。

005：Modular的MAX即时图形编译器

概述

在本教程中，我们将学习Modular公司MAX框架中的图形编译器。这个编译器是深度学习推理平台的核心组件，负责将高级深度学习算子图转换为高效的、可执行的代码。我们将了解其设计哲学、核心架构、扩展机制，以及它如何通过生成Mojo代码来实现高性能和可调试性。

图形编译器高层架构 🏗️

上一节我们概述了图形编译器的目标，本节中我们来看看它的整体架构。

图形编译器接收两个主要输入：

  1. 一个由高级深度学习算子组成的计算图。
  2. 一组用Mojo编写的、参数未绑定的预编译内核包。

其输出是一个符合Modular可执行格式（MEF）的可执行文件，你可以将其理解为专为任意计算图设计的ELF文件。

从方言（Dialect）的角度看，图形编译器主要包含三个处理阶段：

  • RMO/MO阶段：处理模块化算子。RMO是MO的宽松版本，允许更多隐式行为（如自动广播）。
  • MUG阶段：即模块化图生成器，这是进行算子融合等优化的核心阶段。
  • MGP阶段：即模块化粘合原语，它负责对运行时行为以及与设备的交互进行建模。

RMO/MO阶段：高级算子与优化 🔧

上一节我们介绍了编译器的三个阶段，本节我们深入第一个阶段——RMO/MO。

RMO/MO方言定义了约135个核心深度学习算子，例如矩阵乘法（matmul）、激活函数（如relu）和二元运算（如add）。图中的顶点代表这些高级算子，边则代表张量。

该方言原生支持动态形状。例如，在算子属性中，维度N、M、K可以作为符号值存在。为了在静态单赋值形式中维持操作顺序，我们通过chain显式支持副作用。此外，通过可变张量（mutable tensors）的概念支持原地操作。

以下是该阶段典型的优化：

  • 符号折叠：例如，将 add 操作数与零相加的操作折叠掉。
  • 形状推断
  • 常量折叠
  • MLIR提供的通用优化：如公共子表达式消除（CSE）和死代码消除（DCE）。

应对快速演进的行业：可编程性与可扩展性 🛡️

上一节我们了解了基础优化，但深度学习基础设施日新月异。本节我们来看看如何设计一个能适应这种快速变化的编译器。

传统的深度学习编译器设计可以看作一个光谱：

  • 内核中心系统：性能好、易于扩展、可调试、行为可预测，但覆盖率低，任何模型变更都需要人工编写新内核。
  • 编译中心系统：覆盖率高，但通常难以获得峰值性能，编译时间长，行为不可预测，且如果需求超出其编程模型则难以扩展。

Modular的结论是：不存在一个万能的领域特定语言或编程模型能可靠地满足所有不断演进的深度学习推理需求。

因此，我们的解决方案是：通过对可编程性和可扩展性进行投资来规避风险。具体到图形编译器：

  1. 让内核编写成为高效的任务：我们使用专为内核设计的Mojo语言，能够快速实现新硬件的顶尖性能。
  2. 将图形编译器设计为编写内核的基础设施：使其能够支持内核编写。
  3. 为专家让路：硬件专家最了解如何发挥硬件峰值性能，编译器必须允许他们做想做的事。这是图形编译器可扩展性的关键。

核心扩展机制：mod.custom 与 mod.opaque ⚙️

上一节我们提到了为专家让路，本节我们来看看实现这一目标的具体构建模块。

mod.custom 和 mod.opaque 是图形编译器中可扩展性的基石。

  • mod.custom：这是一个高度抽象的操作，可以接受任意数量的操作数和结果，并拥有一个映射到Mojo参数的属性字典。通过它，你可以在图形编译器中做任何你想做的事，只需提供内核实现即可。
  • mod.opaque：类似于 mod.custom 之于内核，mod.opaque 用于处理任意的不透明模块类型。它允许你操作图形编译器本身不理解的数据结构（如非张量类型的缓存），从而实现无需修改编译器本身的扩展。

这种扩展性得益于之前提到的JIT阶段：编译器接收预编译的、参数未绑定的内核包，并在MGP的最后阶段根据输入模型为这些参数提供具体值。

MUG阶段：通过内省实现融合 🔄

上一节我们介绍了扩展机制，本节我们看看如何弥合Mojo内核与融合表示之间的鸿沟。

左边是Mojo中matmul内核的签名，右边是MUG中用于融合的结构化内核表示。如何从前者得到后者？答案是：内省Mojo代码。

Mojo代码被解析后会转换为LIT方言的IR。作为编译器工程师，你可以在Mojo中构建各种抽象（如装饰器、类型），它们会出现在LIT IR中。通过编程方式检查内核签名，我们就可以生成对应的结构化融合表示。

例如，一个包含matmul、bias_add和relu的图，可以在MUG中被融合成一个大的内核。重要的是，这种降低到结构化内核的方式仅基于内核的编写方式，而不关心具体是什么操作。因此，它对于用户通过mod.custom添加的任何自定义操作同样有效。

最终阶段：生成Mojo代码与未来展望 🚀

上一节我们看到了融合如何工作，本节我们来看一切如何最终汇聚为Mojo代码。

图形编译器的最后阶段是生成Mojo代码。输出的Mojo文件会引用用户提供的原始内核。生成的代码几乎是一一对应的，易于理解。此外，编译器还会自动生成运行时样板代码（如异步值的解包），这些代码对用户透明但可供调试。

生成Mojo代码的好处包括：

  • 人类友好：可读、可内省、可调试、可修改。
  • 强大的可能性：任何你能想象在Mojo中实现的优化或模式，理论上都可以通过图形编译器生成。这为未来功能打开了大门。

目前我们已经实现了基础融合（如元素级融合、prologue/epilogue融合）。未来，通过生成Mojo，我们将能够：

  • 添加新的融合类型。
  • 实现Mega Kernel。
  • 改进内核注册机制。
  • 增强垂直调试体验（从Python前端到生成代码）。
  • 优化生成的Mojo代码的可读性。

总结

在本教程中，我们一起学习了Modular MAX图形编译器的核心设计。我们了解了其三层架构（RMO/MO, MUG, MGP），认识了它通过mod.custom和mod.opaque实现的可扩展性设计哲学。我们重点探讨了编译器如何通过内省Mojo内核签名来实现与操作无关的融合优化，并最终生成可读、可调试的Mojo代码。这种强调可编程性、可扩展性并为硬件专家让路的设计，旨在构建一个能够适应深度学习领域快速变化、同时持续提供峰值性能的编译系统。

006：过去、现在与未来 🚀

在本节课中，我们将要学习LLVM中的调度模型。我们将了解它的基本概念、它如何被编译器其他部分使用、它与实际硬件的联系，并探讨其未来的改进方向。

概述 📋

调度模型是LLVM编译器框架中的一个核心组件，它为指令调度器提供了关于指令执行特性的关键信息。理解调度模型对于编写高效的代码生成后端至关重要。

什么是调度模型？ 🤔

调度模型最初是为指令调度而发明的。为了理解它，我们首先需要了解指令调度器的目标。

指令调度器主要有两个任务：第一是提高指令级并行性，第二是减少寄存器溢出。本节课我们主要关注第一个任务。

我们可以通过一个例子来解释。假设我们有一个数据依赖图，箭头表示从数据生产者到消费者的依赖关系。我们的目标是减少流水线停顿。

例如，一个加载指令通常有加载延迟。如果我们什么都不做，流水线就会在加载指令完成前浪费很多周期。我们希望做的是，在加载指令和它的消费者（比如一个加法指令）之间，填充一些无关的指令。

另一个例子是，如果有三条可以独立调度的指令，但它们竞争相同的硬件资源，我们就需要知道如何安排它们以避免冲突。

指令调度需要的关键信息 🔑

从上面的例子，我们可以总结出指令调度需要从调度模型获取的三个核心要素：

  1. 延迟：指令执行完成所需的总周期数。
  2. 资源：指令执行时需要使用哪些硬件执行单元。
  3. 占用时间：指令占用特定资源的时间长度。

一个最简单的调度模型可以是一个巨大的表格，其中列出了每条指令的延迟、使用的资源和占用时间。

从硬件视角看调度模型 ⚙️

上一节我们介绍了指令调度需要的信息，本节中我们来看看这些概念如何映射到实际的硬件。

最初的调度模型称为“行程模型”。它将指令建模为多个阶段，每个阶段都定义了占用时间和使用的资源。整个指令的执行时间就是所有阶段时间的总和。

这很容易让人联想到传统的处理器流水线模型。延迟对应指令流经各个流水线阶段的总时间。占用时间特指指令在执行阶段花费的时间。

当我们引入超标量架构时，情况变得更有趣。超标量意味着复制多个相同的执行流水线（称为执行单元），以便同时执行指令。

  • 资源：现在，“资源”具体指代哪个执行单元。例如，指令A可能只能在单元0上运行，而指令B可以在单元0或单元1上运行。
  • 占用时间：含义保持不变，指指令占用特定执行单元的时间。

现代的调度模型：基于读写操作 🆕

现在，我们来看看当前一代的调度模型。它不再将指令拆分为多个阶段，而是从单个指令出发，专注于其操作数。

它将操作数分为“写”和“读”。对于“写”操作，会分配一个令牌，并用这个令牌映射到一个称为 WriteRes 的实体。这个WriteRes包含了我们之前提到的三个要素：使用的资源、延迟和占用时间。

占用时间由两个字段表示：AcquireAtCycle 和 ReleaseAtCycle。我们稍后会详细解释。

以下是一个定义指令及其WriteRes的示例代码：

def : InstRW<[MyWriteRes], (instrs MY_INSTR)>;

def MyWriteRes : WriteRes<[PipeA, PipeB], [1, 33]>;
  • MyWriteRes 定义了该指令使用 PipeA 和 PipeB 资源。
  • [1, 33] 表示延迟。通常，延迟值等于最长的占用时间。

资源的表示

资源通过 ProcResource 声明。有时，一条指令可以在多个资源上执行，这时使用 ProcResourceGroup。如果一条指令需要同时使用多个资源，则在WriteRes记录中列出所有资源。

占用时间的表示

占用时间通过 AcquireAtCycle 和 ReleaseAtCycle 这两个数组来表达。数组的元素数量等于该指令可使用的资源种类数。

例如，对于使用两种资源的指令，数组有两个元素。第一个元素（比如对应PipeB）的 AcquireAtCycle 是0，ReleaseAtCycle 是1，表示它在PipeB上占用1个周期。第二个元素（对应PipeA）的 AcquireAtCycle 是0，ReleaseAtCycle 是33，表示它在PipeA上占用33个周期。

AcquireAtCycle 不一定从0开始，它可以是非零值，用于表示更复杂的流水线行为。

调度模型如何指导调度？ 🧭

将WriteRes信息整合后，调度器就可以工作了。

  • 防止数据冒险：对于加载指令，调度器知道其延迟，因此会将依赖其结果的加法指令适当推后执行，避免加法指令空等数据。
  • 防止结构冒险：对于竞争相同资源的两条乘法指令，调度器知道第二条必须等待第一条释放资源后才能开始。而对于使用不同资源的乘法和减法指令，它们可以被安排在同一时间执行。

硬件也有自己的机制来避免这些冒险，即乱序执行。

乱序执行与调度模型 🔄

乱序执行的核心是在执行单元前增加一个缓冲区（重排序缓冲区）。指令先被分派到缓冲区，等到其所有数据依赖和资源冲突解决后，才被发射到执行单元。

工具 llvm-mca 可以模拟这种行为，它同样使用调度模型来预测指令在CPU上的执行时间线。在模拟中，我们可以看到，没有数据依赖的指令即使代码顺序在后，也可能先开始执行。

对于结构冒险，如果多条乘法指令只能在一个特定单元上执行，而加法指令可以在任何单元上执行，那么加法指令可能先于后面的乘法指令开始执行，因为它们不竞争资源。

调度模型中的缓冲区建模 📊

为了对乱序执行进行建模，调度模型需要描述缓冲区。

  • 统一重排序缓冲区：所有流水线共享一个大的缓冲区。在模型中，这通过设置 BufferSize = -1 来暗示，并使用 MicroOpBufferSize 表示缓冲区总大小。
  • 分离的重排序缓冲区：每个执行单元有自己的缓冲区，通过 BufferSize 属性指定大小。
  • 共享缓冲区：多个执行单元共享一个缓冲区，这提供了更大的调度灵活性。统一缓冲区可以看作是这种模式的最通用形式。

对于顺序执行核心，通常设置 BufferSize = 0，表示没有缓冲区。但有一种特殊的“顺序核心”会将 BufferSize 设为1，表示有一个微小的暂存区。虽然两者在硬件上都是顺序执行，但在调度模型中的表示不同，这引出了下一个关键点。

不同使用者对调度模型的解读 👀

调度模型本身只是提供数据，关键在于不同的使用者如何解读它。llvm-mca 和机器指令调度器（编译器的一部分）对模型的解读就有显著不同。

最大的区别在于对缓冲区大小的解读：

  • llvm-mca 的解读与硬件一致，用于模拟指令在缓冲区中的行为。
  • 机器调度器则完全不使用缓冲区大小信息。因为它无法在编译时预测硬件调度器在运行时的具体行为。

对于顺序调度，机器调度器会尝试在编译时预测停顿。它使用内部数据结构来模拟指令流，并严格检查数据冒险和结构冒险。只有当前没有冒险的指令才会被考虑调度。

而对于乱序核心（或 BufferSize=0 的顺序核心），机器调度器则乐观得多。它假设硬件会处理好冒险，因此几乎不设限制地允许指令进入可调度队列。

这种差异可能导致机器调度器对某些硬件过于乐观，期望硬件具备其实际没有的能力（比如对 BufferSize=0 的核心进行乱序调度）。但有时，这种乐观也有用处，例如，它可能允许一条会带来轻微停顿但能极大降低寄存器压力的指令被调度，从而避免更严重的性能损失。

调度模型的其他潜在用途 💡

除了指令调度和 llvm-mca，调度模型还可以在其他地方发挥作用：

以下是当前已经使用或未来可能使用调度模型的一些场景：

  1. 机器轨迹度量：这是一个未被充分利用的强大工具。它利用块概率信息，估算代码路径的延迟和资源使用情况。目前仅用于机器指令组合器和早期If转换优化。
  2. 机器流水线化器：用于软件流水线优化，它使用所有调度模型信息来重叠循环不同迭代的指令。
  3. TTI成本模型：TargetTransformInfo 目前为LLVM IR指令提供代价信息。未来可以考虑让其基于调度模型来提供延迟等数据，实现“单一事实来源”。挑战在于LLVM IR指令和机器指令之间存在巨大差异。
  4. 指导循环展开：利用调度模型信息来指导循环展开策略，以改善展开后的指令调度。
  5. 验证调度结果：在机器调度器完成调度后，使用 llvm-mca 来验证调度结果的质量。
  6. 改进指令选择模式排序：目前指令选择模式的顺序可能影响调度结果。一个存在了20年的想法是利用调度模型中的延迟信息来对这些模式进行排序，可能带来性能提升。

总结 🎯

本节课中我们一起学习了LLVM调度模型的核心内容。

我们了解到调度模型为指令调度提供了延迟、资源和占用时间这三个关键信息。我们看到了这些概念如何映射到处理器的流水线和执行单元。更重要的是，我们认识到调度模型的精髓不在于模型本身的数据，而在于不同的工具（如编译器调度器和性能分析工具）如何解读和运用这些数据。

最后，我们探讨了调度模型当前的应用以及未来可能的改进方向，例如统一代价模型、验证调度结果和优化指令选择等。理解这些将帮助我们更好地利用调度模型来生成高性能的代码。

007：核心设计与实现

在本教程中，我们将探讨为 CUTLASS 实现的 Python 领域特定语言（DSL）及其底层基础设施。我们将了解其设计动机、核心架构、编译流程以及如何平衡 Python 的灵活性与高性能计算的需求。

设计动机与目标

大家好，我是郭藏，在英伟达担任编译器工程师。今天我们将讨论我们为 CUTLASS 实现的 Python DSL 及其基础设施。

你可能会问，CUTLASS 的 Python DSL 是什么？CUTLASS 不是一个 C++ 库吗？是的，它曾经是一个库。但当这个项目变得非常重要时，我们为其实现了一个编译器，并在其之上构建了 DSL。我们的 DSL 称为 CUTLASS DSL（CUTLASS DSL），它是一个基于装饰器的 DSL。你基本上可以通过 pip install 安装并立即开始使用。我们为主机和设备代码提供了不同的装饰器。

由于这是英伟达 GPU 和 CUTLASS，你可能会问它的速度如何，加速效果怎样？你可以获得非常快的速度，因为这种 DSL 的设计理念是为你提供完全的控制和抽象。通过使用这些功能，你确实可以获得峰值性能。我们已经为不同架构提供了许多示例内核，并提供了像 Jupyter Notebook 这样的教程。我相信你会找到大量材料来理解如何使用这个 DSL。

但今天的讨论重点不在于此，因为这是一个编译器会议。我想告诉你我们为什么实现这个 DSL，以及在实现过程中做了哪些设计选择。我不会展示任何性能数据，请不要失望。

核心架构：基于 MLIR 的编译栈

那么，让我们开始吧。这是我们的起点，即我们的蓝图。我们认为 MLIR 层是实现抽象的好地方，因为我们可以实现方言。我们从 MLIR Affine 和 MLIR SCF 方言开始。

然后我们对其进行了扩展。扩展意味着我们将 CUTLASS 代数建模为 MLIR 中的一个方言，位于 MLIR Affine 和 SCF 方言之上。从这个方言出发，我们实现了转换，最终生成 MLIR NVVM 方言。因此，我们的关键方言是 NVVM 方言。我们实际上正在朝这个方向推进。这是一个完整的方言，不再有内联汇编。我们直接生成 LLVM IR。之后，我们当然会转到 LLVM 方言，然后生成 PTX。

这就是我们当前的工作流程。现在我们希望使其可编程。为了实现这一点，我希望从 Python 中交付这个功能。当我想要这样做时，我看到了 Python 语言的庞大性。它非常庞大，可能超出了子集的范围，而且我没有任何工具可以将任何 Python 语言映射到我的 MLIR，因为我们还没有构建这样的工具。

那么，我的问题是：如何将 MLIR 捕获到 Python 中？如何在 Python 中对 MLIR 进行建模？实际上，这甚至不是我的问题，因为我不是在开发一门新语言，这不是我们的目标。我们的目标是交付用于 CUTLASS 的 DSL。而 CUTLASS 关乎性能。所以我的问题是：哪些 Python 特性能够让我快速编写 GPU 内核，同时又不牺牲性能？我不能为了使用所有 Python 特性而牺牲性能。这就是我们的目标。

DSL 实现方法对比

在开始之前，让我先谈谈 DSL 的实现方式。实现 DSL 有多种方法，我将讨论其中的三种。

第一种是基于 Python 抽象语法树（AST） 的方法。你基本上获取 Python AST，然后将其映射到你的 MLIR 或类似的中介表示（IR）。

第二种是基于 字节码 的方法。你获取 Python 字节码，然后生成 IR。如果你来自 C++ 或 Rust/Fortran 背景，这通常被称为编译器，是大家常用的方法。

第三种是 追踪（Tracing） 方法，它根本上是不同的。你不处理抽象语法树，而是运行程序并在执行过程中捕获信息。这是非常不同的。

让我们深入探讨一下。我想向你展示基于 AST 的 DSL 思路：你有一个 Python 装饰函数，通过解释器生成 AST，然后将其映射到 IR。这就是整个想法。

但是，让我们看一个更现实的例子。这里我展示了一个例子。它非常简单，但在机器学习中非常典型和复杂。我们所做的就是执行一个通用矩阵乘法（GEMM）。在 GEMM 之后，我们有一个收尾操作（epilogue），例如，我对结果进行一些处理。

在这个例子中，对于收尾操作，我有两个类。一个类基本上什么都不做，但另一个类执行 ReLU 激活函数。我使用了多态风格，因为它们是很好的概念，使我的代码可读性强。

现在，让我们为这个函数生成 AST。这是为装饰函数生成的 AST（不是所有代码的 AST，因为我只对实现基于装饰器的 DSL 的装饰函数感兴趣）。

这是我的 AST。让我们逐行理解这些节点。

第一行是 for 循环。for 循环存在于 AST 中。这很好，因为我不需要理解循环周期和做其他事情。它就在那里，非常酷。所以它对我有用。

对于 GEMM，我有一个函数调用。为简单起见，我们假设 GEMM 已经在一个函数中实现，并且速度非常快。这对于 AST 来说没问题。

对于收尾操作，糟糕的事情发生了，因为 Python 是动态类型的，我无法在 AST 层面理解这个 epilogue 是什么。我在调用什么并不清楚。我需要做的是插入更多的装饰器，为每个 epilogue 相关的东西进行分析。这非常困难。

让我们总结一下基于 AST 方法的优缺点。

优点：你的程序结构清晰。循环、控制流都在那里，一切都很清晰，这很好。

缺点：你无法免费获得 Python 的特性。你必须在你自己的 Python AST 中实现所有你想要的功能。当你这样做时，我确信它将不再是 Python，而是变成了别的东西，变成了 DSL。人们将不得不在你的 DSL 中思考，而不是在 Python 中思考。而且你的 DSL 发展会很慢，因为你必须先实现功能，然后才能映射到 DSL，这需要时间。所以这可能不是最好的主意。

现在，让我们转到第二种方法，即基于字节码的方法，因为它与 AST 方法非常相似，只是层次更低。但第三种方法是根本不同的，我想解释一下。

在这里，我们不处理 AST，而是用解释器执行程序。让我们执行一下。由于是解释执行，我也会展示解释器在做什么。

首先，它读取所有代码，如果遇到内核函数调用，我们的装饰器就会介入。在右侧，我展示了生成的 MLIR 代码，因为我们在执行过程中生成代码。

对于你的函数，我们生成某种 MLIR 函数。这很酷。然后对于 for 循环，我们执行它，但糟糕的事情又发生了。现在 Python 解释器执行了我的循环，因为在这种情况下，Python 正在执行 Python 代码，而我正在捕获我所捕获的内容。for 循环是 Python 的，所以我的循环消失了。如果我有一个控制流，它也会消失。所以这并不好。

但是，当我们运行 epilogue 时，神奇的事情发生了，发生了很多很多事情。这里发生的是：我构建了类，构造了类，并且有一个多态函数调用，它起作用了。此外，函数调用是内联的。所以我生成的代码非常简洁漂亮。我没有为此付出任何代价。一切都生成得很好，我免费获得了 Python 的特性，因为 Python 解释器比任何人都更懂 Python。我免费获得了这些特性。

现在让我们转到第二个内核的调用。在这种情况下，我有一个不同的激活函数。我将保留所有步骤，但直接看最后一步。在最后一步，你会看到一切照常运行。我们重新生成了 IR，但 epilogue 部分不同了。现在我可以使用 ReLU 激活函数了。

那么，我们也来总结一下这种方法。

优点：很酷，因为我可以免费使用 Python 的特性，并且可以使用 Python 进行元编程。

缺点：解释器运行得太自由了，我无法控制它，它吞噬了我的程序控制流和所有函数调用。是的，这可能也不是最好的主意。

那么，我需要选择哪一种呢？我卡在这里了，但我并没有，因为在英伟达，我可以询问编写内核的人他们需要什么。我收集了两个事实。

第一个事实非常明显：快速的 GPU 内核算法非常简单。没有函数调用，没有面向对象编程，没有多态，没有复杂的分支或深度嵌套的逻辑。你可以在 GPU 上拥有所有这些，但不是在快速内核中，因为我们追求的是峰值性能。

但是，当我们查看内核代码时，它们使用了元编程等概念，比如使用类，因为这很好。它使我们的代码可读性强，就像 C++ 一样，同时也保持了代码的整洁，提高了可读性。我们喜欢这些东西，但我们不想为此付出性能代价。

我们的解决方案：结合 AST 与追踪

因此，我们的想法是结合这两种方法，即结合 AST 和追踪。我们让用户与 Python 解释器成为朋友。解释器不会随意运行，你可以阻止它，这意味着捕获 IR。如果你想让 Python 运行，那么你可以让它运行。我们如何做到这一点呢？我们引入了 常量表达式 的概念。如果你有常量表达式，我们将在 Python 解释器上运行它。

那么，什么可以是常量表达式呢？函数参数可以是，变量可以是，控制流可以是，形状（如果你关心的话）也可以是常量表达式。但由你选择，而不是我，因为你知道什么必须被捕获，什么不需要被捕获。

那么，我们的编译流程是怎样的呢？我们有两个阶段和一个预阶段。让我从头开始解释。

从 Python 代码开始，我们进入预阶段。在这个预阶段，我们读取你的 Python AST，并生成另一个 Python 代码，这有点像元 Python。在这里，我们理解你的控制流和程序结构，并生成回调函数。然后你进入元阶段。在这个元阶段，我们有了元 Python 程序，并用 Python 解释器执行它。在这个过程中，我们进行常量折叠、实现回调、进行追踪、执行部分求值、类型推断等，所有事情都在这里完成。然后，这个阶段产生一个对象，比如 MLIR（MLIR 是静态类型的），这个对象基本上就是 CUTLASS 方言。接着我进入 MLIR 对象阶段。在这里，我们最终确定代码，生成二进制文件。

让我们举一些例子来理解如何在元阶段调试或理解发生了什么。调试工具并不令人惊讶，就是 print，因为 print 在追踪阶段（即我们的元阶段）工作。

我将展示三个例子。让我们从头开始。

第一个例子：我们有 a 和 b，它们不是常量表达式，它们是动态的。我有一个结果 c，并用它们进行计算。当我打印 c 时，我看不到数据结果，因为 c 在元阶段没有被求值，它没有在 Python 解释器中求值。你必须在目标架构上执行这个程序才能看到结果。

第二个例子：我们再次有 a 和 b，但 b 是一个常量表达式（因为你可能这样标记它）。a 是一个 Python 值，因此根据定义也是常量表达式。当你进行求值时，你立即在打印中看到它的值，因为我们在元阶段用 Python 解释器对它进行了求值。所以你的结果 c 有其值。

第三个例子是混合的：其中一个 b 是常量，a 是动态的。在这种情况下，我们显然会将常量提升为动态，你看不到结果。

当然，我们也有常规的 printf 或 CUTLASS 的 printf。如果你使用这个 printf，它将在二进制文件执行后打印结果，可能来自你的 GPU。所以这是不同的。

CUTLASS DSL 应用示例

现在，让我们用 CUTLASS DSL 来写我之前展示的例子。我们实际上不会改变太多东西。但我们在这个程序中真正想要的是两件事：我想使用 epilogue，并且我希望这个 epilogue 在 Python 解释器中运行，因为我使用了一些元编程；对于 for 循环，我实际上想捕获它。这并不难。

对于 epilogue，我将其注释为常量表达式，因为我想使用 Python 解释器。对于 for 循环，在循环中，我不对其进行任何注释，因为我想捕获这个值。如果我们不注释任何东西，并且使用 Python 的常规 range，你会看到循环变量 i。目前，我们使用 SCF 方言来处理它。

但如果你想将其标记为常量表达式，也可以。在这种情况下，你的循环将由 Python 运行，因此会自动展开。

对于 epilogue，是的，由于它是常量，我们将由 Python 解释器进行求值。

这就是我们的 DSL 代码生成部分的基本工作原理。

扩展性与多方言支持

现在我想谈谈这个 DSL 如何与英伟达内部的其他方言协同工作，因为在英伟达，我们内部有多个方言，我们希望使用相同的 DSL 来编程这些方言或其他方言。

当我们查看 DSL 的特性时，我在这里展示的特性实际上是目标无关的。我的意思是，当你从一个目标转到另一个目标时，它们并没有太大差异。例如，AOT 编译在某种程度上是目标相关的，或者像 GPU 缓存这样的特性是目标无关的。或者你需要运行 MLIR 执行引擎和即时编译，所有这些事情，比如 LLVM，都是目标无关的。

有两件大事不是目标无关的，它们是类型推断和由 Python AST 生成的回调函数。

但我们确保它们是可扩展的。

例如，让我们看看类型推断。我们有内置类型，这些内置类型被每个其他方言共享，因为每个人都使用标量类型。

但是当你有自己的 DSL 实现时（在这种情况下，我们讨论的是 CUTLASS DSL），你将拥有自己的类型，这些类型可以与现有的 DSL 类型一起工作。所以我们的类型推断是可扩展的。

另一个是生成回调函数。之前我展示了编译流程，并提到我们有 Python 代码，然后进入预阶段，在那里我们读取 AST 并生成元代码（元 Python 代码）。在这种情况下，我们的元 Python 代码会生成回调函数。你必须为你自己的 DSL 实现这些回调函数。例如，对于 CUTLASS DSL，我们实现了这些回调函数来生成 SCF 方言以处理控制流。对于其他事情，你可以实现任何其他功能。例如，如果你想实现另一个 DSL（假设叫 my_dsl），你可以实现这些回调函数，生成存在于你方言中的任何其他控制流，这完全没问题。因为我们希望确保这个元 Python 代码是相同的，并且可以与我们希望使用此 DSL 基础设施的任何类型的方言一起工作。

最后我想谈谈调用约定，因为你可能会问：你的 DSL 到底有多 Python？我们确保你可以从任何地方调用 Python。这是我们的调用约定。我提到我们有两个装饰器：一个装饰器是内核（GPU 函数），另一个装饰器是设备函数（可以是主机或设备函数）。正如你在表格中看到的，我们基本上可以从任何地方调用任何地方。这里缺少两件事：例如，最后一个“内核调用内核”我们不支持，我们没有找到任何有用的用例，所以不支持。另一个是 Python 函数直接调用内核，我们也不支持，同样因为没有找到有用的用例，因为我们总是需要经过主机来生成线程束脚本，然后再回到设备。但我的意思是，如果你找到有用的用例，我们可以实现这两个功能。所以这并不是根本性的困难。

总结

好了，让我总结一下我的演讲。

我们已经实现了这个基础设施，并在 CUTLASS 仓库下交付了 DSL CUTLASS DSL。它是目标无关的，可以与许多方言协同工作，目前内部已与多个方言一起工作。

我们的 DSL 分为两个阶段和一个预阶段。在预阶段，我们读取 Python AST 以捕获程序结构。在元阶段，我们运行 Python 解释器。在对象阶段，我们使用 MLIR 编译静态类型代码。

谢谢大家，我的演讲到此结束。

本节课中我们一起学习了 CUTLASS Python DSL 基础设施的核心设计。我们探讨了其结合 AST 分析与运行时追踪的独特编译流程，理解了常量表达式在平衡 Python 灵活性与生成高性能 GPU 代码之间的关键作用，并了解了该 DSL 如何设计为可扩展、目标无关的系统，以支持多种 MLIR 方言。

008：P08-P07_针对LLVM的RISC-V后端的翻译验证

在本教程中，我们将学习如何使用翻译验证工具来检查LLVM的RISC-V后端是否存在编译错误。我们将介绍工具的工作原理、具体步骤，并通过一个实例来理解如何发现和修复后端优化中的问题。

概述：什么是翻译验证？🔍

翻译验证是一种形式化验证方法，用于确保编译器优化或代码转换的正确性。具体来说，它验证转换后的目标代码是否是源代码的一个“精化”。这意味着目标代码的行为必须与源代码一致，或者更精确（例如，不会引入未定义行为）。

LLVM社区有一个名为Alive2的工具，它使用Z3求解器来验证中间表示（IR）层面的优化。然而，我们的重点是将这种方法应用于后端，特别是针对RISC-V架构的机器码生成。

翻译验证的工作流程🔄

上一节我们介绍了翻译验证的概念，本节中我们来看看针对RISC-V后端的具体验证流程是如何分步进行的。

整个验证过程可以概括为四个主要步骤：

  1. 调用LLVM后端：输入LLVM IR，让后端生成RISC-V汇编代码。
  2. 将汇编代码提升回IR：使用一个专门的“提升器”，将生成的RISC-V汇编代码转换回LLVM IR形式。这一步是关键，它创建了一个与原始IR可比较的中间表示。
  3. 优化提升后的IR：使用opt工具对提升得到的IR进行简化，去除不必要的细节（如寄存器分配、栈操作），得到一个更干净、更核心的IR用于比较。
  4. 进行精化验证：使用Alive2工具验证第3步得到的目标IR是否是第1步原始源IR的一个精化。

如果验证通过，说明整个后端转换过程（IR -> 汇编 -> 提升后的IR）是保持语义正确的。如果Alive2报告错误，则意味着其中某一步转换不是精化，可能存在编译错误。

RISC-V后端验证工具：RISC-V TV🛠️

我们基于上述流程开发了RISC-V TV工具。它与之前用于ARM64后端的RTV工具思路类似，但主要区别在于“提升器”的实现。

RISC-V的提升器实现起来相对简单，主要原因是RISC-V指令集架构没有状态标志位（如x86的EFLAGS）。因此，我们不需要像处理x86那样去模拟和撤销标志位的计算，这大大简化了提升过程。

接下来，我们深入了解一下提升器是如何工作的。

提升器的工作原理🧩

提升器的核心任务是模拟一个RISC-V函数的执行环境，并将每条机器指令映射到语义等价的LLVM IR操作。以下是其工作步骤：

  1. 分配存储空间：首先，在内存中分配空间来模拟RISC-V的寄存器文件。同时，分配一个栈空间用于处理函数调用约定（ABI）。
  2. 翻译每条指令：对于汇编代码中的每条RISC-V指令：
    • 从模拟的寄存器文件中加载（load）操作数。
    • 使用语义等价的LLVM IR指令执行操作。
    • 将结果存储（store）回目标寄存器。

以下是一个简单的例子，展示了提升器内部如何处理add指令：

; 假设 %reg_file 是模拟寄存器文件的基础指针
%val1 = load i32, i32* %reg_file_offset_1 ; 加载第一个操作数
%val2 = load i32, i32* %reg_file_offset_2 ; 加载第二个操作数
%sum = add i32 %val1, %val2                ; 执行LLVM的add操作
store i32 %sum, i32* %reg_file_offset_dest ; 存回目标寄存器

这个例子很简单，因为LLVM的add指令语义与RISC-V的add指令完全相同。但并非所有指令都如此。

有时会遇到边缘情况，或者某些RISC-V指令在LLVM中没有直接对应的IR操作。这时，我们需要用一系列LLVM指令来完整实现该指令的语义。在这个过程中，我们的目标不是生成高性能代码，而是生成对求解器（如Z3）“友好”的、易于验证的IR序列。通常，这意味着一连串清晰的位操作指令，而不是循环或复杂的向量操作。

提升后代码的结构📊

经过提升，我们会得到大量的LLVM IR代码。其中大部分代码（例如黄色的寄存器分配、栈操作、ABI包装等）是用于构建模拟环境的样板代码。如果我们过滤掉这些，就能看到核心的指令映射。

最终得到的目标函数结构清晰：

  • 序言：分配和初始化模拟的寄存器与栈。
  • 主体：每个RISC-V指令通常对应一个基本块，其中包含加载操作数、执行操作、存储结果。
  • 尾声：处理ABI要求，如位宽扩展（零扩展或符号扩展），然后返回。

有了这个系统，我们就可以开始用测试用例来寻找后端bug了。

测试用例的来源与模糊测试🔬

为了有效地发现bug，我们需要大量的、多样化的测试输入。我们主要从以下几个来源获取测试：

以下是我们的测试用例来源：

  • LLVM测试套件：包含大量专门设计用于测试后端模式的用例。
  • YARPGen：生成随机的C函数，编译成IR后作为输入。
  • Yuf：对现有的IR进行细微变异，例如插入指令、修改参数或改变位宽，以探索边缘情况。
  • Elegant Mutation-Based Bug Seeding：另一个有效的变异生成器，帮助我们发现了许多问题。

实例分析：发现并修复一个Bug🐛

让我们通过一个具体的例子来看看如何发现和修复bug。考虑以下LLVM IR代码，它使用了cttz（计数尾随零） intrinsic：

define i8 @example(i8 %x) {
  %iszero = icmp eq i8 %x, 0
  %cttz = call i8 @llvm.cttz.i8(i8 %x, i1 false) ; false 表示输入为0时返回poison
  %result = select i1 %iszero, i8 0, i8 %cttz
  ret i8 %result
}

这个模式的意图是：如果输入是0，则返回0；否则返回尾随零的个数。然而，RISC-V的ctz指令在输入为0时的行为是返回数据的位宽（对于i8就是8），而不是LLVM IR中poison或我们期望的0。

在早期的RISC-V后端优化中，编译器试图用一种巧妙的方式处理这个差异。它生成的汇编可能包含设置标志位、进行掩码操作来模拟select行为。但是，当模糊测试工具将i8类型变异为i7（非2的幂次方位宽）时，问题暴露了：用于计算的掩码值错误，导致一半的输入结果出错。

根本原因是，掩码的计算方式（位宽 - 1）只对2的幂次方位宽有效。我们随后提交了一个PR，修改了后端的代码生成逻辑，现在这个优化对于非2的幂次方的类型也能正确工作了。

成果与总结📈

通过大规模的翻译验证和模糊测试，我们取得了以下成果：

  • 对于ARM64后端（RTV项目）：发现了45个此前未知的编译错误。
  • 对于RISC-V后端（RISC-V TV）：目前只发现了1个RISC-V特有的编译错误。

这个数量差异可能令人惊讶，但有几个原因：

  1. 许多后端代码是共享的。在ARM64后端上通过模糊测试发现的bug，在被修复后，也惠及了RISC-V等其他后端。
  2. RISC-V社区（如Craig Topper等开发者）已经积极修复了许多已知问题。
  3. 我们的测试可能尚未覆盖所有RISC-V扩展（例如V向量扩展）。

这项工作告诉我们，对一个后端进行深入的翻译验证和模糊测试，由于代码共享，实际上对所有LLVM支持的后端都有益处。它不仅帮助发现bug，也使得翻译验证本身变得更容易实施。

未来方向🚀

本节课中我们一起学习了针对LLVM RISC-V后端的翻译验证方法。展望未来，可能的工作方向包括：

  • 支持更多的RISC-V ISA扩展（如V扩展）进行测试。
  • 改进模糊测试的变异策略，以生成更有效、更复杂的测试用例。
  • 进一步分析当前方法可能遗漏的bug类别。

009：Mojo的GPU编译

在本节课中，我们将要学习Mojo语言如何实现GPU编译。Mojo是一种Python风格的系统编程语言，它结合了Python的易用性与C++级别的性能，并提供了对CPU和GPU的统一编程模型。

概述

Mojo是Modular公司正在构建的一种Python风格的系统编程语言。它看起来像Python，但拥有支持泛型编程、强类型系统和内存安全等强大特性。我们认为它是将Python扩展到CPU和GPU的最佳方式。Mojo非常快，具备类似C/C++的性能。它也是Modular Max推理引擎的基础之一。Mojo提供了一个统一的编程模型，可以同时用于CPU和GPU，从而释放标准CUDA和ROCm的全部能力，但用户无需直接使用这些供应商工具链。

Mojo的GPU编程模型

上一节我们介绍了Mojo的基本概念，本节中我们来看看Mojo如何实现GPU编程。

Mojo的大部分GPU功能实际上是Mojo标准库的一部分，这意味着它们是Mojo代码，而非编译器魔法，这也使得编译过程更加快速。

以下是一个统一的Mojo程序示例，其中定义了GPU内核函数和CPU驱动代码：

# 定义一个GPU内核函数
fn kernel_function(...):
    # GPU内核代码

# CPU驱动代码
fn main():
    # 获取设备上下文
    context = get_device_context()
    # 在设备上创建缓冲区
    buffer = create_buffer_on_device(context)
    # 调用内核编译和启动函数
    result = inQ_function(kernel_function, buffer)
    # 将缓冲区数据带回主机
    host_data = bring_buffer_back_to_host(buffer)

在这个模型中，程序开头定义了GPU内核，CPU驱动代码负责获取设备上下文、在设备上创建缓冲区，然后通过inQ_function调用内核函数。编译器在底层会编译并启动该内核，之后可以将缓冲区数据带回主机。

GPU编译流程

了解了编程模型后，我们深入探讨Mojo的GPU编译流程。Mojo编译包含多个主要阶段。

Mojo源代码首先经过解析器，进入参数化中间表示（IR）领域，在此进行语义检查和初步优化。随后进入实例化阶段，类似于C++模板实例化。在此之后，所有参数都会被具体化，然后进行更多优化，最终生成二进制代码。

GPU编译发生在实例化阶段。这是因为，如前一张幻灯片所示，GPU内核函数是作为参数提供给inQ_function进行编译的。

在实例化之前，GPU代码和主机代码都位于同一个MLIR模块中。在实例化期间，当我们知道需要将某些内核编译到NVIDIA后端时，我们会将这些内核切片到一个单独的MLIR模块中，并将目标信息（此处是NVIDIA）附加到该模块。然后，我们会运行与CPU端相同的完整编译流程，只是目标信息不同。接着，我们使用并行化后端将不同的内核切片到单独的LLVM模块中，进而生成内核代码。

在此过程之后，假设我们尝试编译两个内核，我们将为每个内核获得一个PTX（并行线程执行）文件。然后，我们将这个编译结果作为KGn_compile_offload函数的返回值插入，以便主机端可以消费它，无论是用于打印还是启动内核运行。

编译控制与调试支持

Mojo的编译流程还支持对编译过程的控制，以方便调试。

我们可以控制希望在哪个阶段停止内核的编译。例如，我们可以选择查看MLIR中间表示、汇编IR或最终的目标代码。这一切都可以在库级别进行控制，库会告诉编译器何时停止。

总结与优势

本节课中我们一起学习了Mojo的GPU编译机制。作为总结，Mojo提供了一种统一的方式，使用同一种语言编写CPU和GPU内核。其大部分功能由Mojo库驱动，这使得编译器更加简单，无需进行大量复杂的魔法操作来适应编程模型。

MLIR在这里对我们非常有帮助，它实现了无缝的编译器集成，我们可以通过添加不同的操作来扩展编译器的能力，而无需改变语言本身。我们还可以直接使用上游的方言（如NVVM、ROCm DL和LLVM），并将它们作为Mojo方言插入到库中，因为Mojo也是MLIR方言的语法糖。

这种方法不仅限于GPU。我们相信它普遍适用于任何其他加速器，只要它们有一个后端，其流程将与上述流程非常相似。或者，如果它基于MLIR，我们可以将其作为Mojo的领域特定语言（DSL）插入以提供支持。

目前，所有的Mojo GPU内核都是开源的。你可以在Mojo代码中查看我们如何实现这些高性能内核。我们将在明天的会后讨论环节进行交流，如果你想了解更多，欢迎前来与我们交谈。

谢谢大家。

010：LLVM中的时序抗干扰编码支持

概述

在本节课中，我们将要学习编译器如何无意中破坏密码学代码的安全性，并介绍LLVM中引入的一种新内建函数（builtin intrinsic），它能为密码学操作提供跨优化级别和架构的恒定时间保证，从而抵御时序侧信道攻击。

编译器优化：一把双刃剑 🛡️

编译器在优化代码方面表现得极其出色，甚至可能过于出色。近年来，其优化能力持续增强，涌现出许多广为人知的特性，例如：

  • 循环向量化
  • 死代码消除
  • 常量传播

对于大多数代码库而言，这些优化特性如同魔法。然而，对于密码学库或密码学算法而言，它们却可能成为一种负担，并从根本上破坏其安全性。

一个破坏性的例子 🔓

密码学家编写的代码是经过精心设计的位操作代码，期望其执行时间是恒定的（constant time）。这段代码进入LLVM编译流水线，经过中端和后端的多次优化传递后，生成了如下的x86汇编代码。

; 示例：存在数据依赖分支的汇编
cmp    %rax, %rbx
je     .Lsecret_block    ; 条件跳转指令

正如你所见，这里有一条JE（Jump if Equal）指令，这实质上意味着代码中存在数据依赖的分支。这为时序侧信道攻击打开了大门。攻击者能够测量不同指令执行时间的微小差异。

具体来说，如果变量i等于秘密索引secret_index，程序就会发生跳转。这次跳转比不跳转需要多几个CPU周期。正因为如此，攻击者本质上能够推断或泄露秘密信息。

问题的严重性 ⚠️

虽然这看起来是一个相对简单的优化问题，但它却导致了影响数百万乃至数十亿系统的灾难性漏洞。

我们可以更进一步了解。苏黎世联邦理工学院的一个研究小组进行了一项名为“Breaking VAd”的研究。他们编译了八个密码学库，针对不同版本的LLVM和GCC编译器、不同的优化级别和不同的处理器架构进行了测试。结果显示，在超过44,000个测试配置中，存在编译器引入的漏洞，影响了诸如BoringSSL、OpenSSL、HHL等主流密码学库。

这从另一个角度说明，存在大规模漏洞快速影响整个生态系统的可能性。

现有的缓解措施及其不足 🚧

业界已有一些尝试缓解此问题的工作，但均未能成功并入LLVM主线，因此我们从未能从中受益。

你可能会想，密码学家如今是如何应对的呢？他们的解决方案并不理想：

  • 使用内联汇编：这种方法不可移植，也难以维护。
  • 采用位操作技巧：希望借此绕过优化，但已被证明并不可靠，有时仍会被优化折叠，导致时序侧信道漏洞。
  • 最终手段：直接禁用优化，但这会严重牺牲性能。

解决方案：__builtin_constant_time 内建函数 🛠️

那么，我们该如何解决呢？我们在Chlippis的团队创建了__builtin_constant_time内建函数。它能在所有优化级别和多种架构上提供恒定时间保证。

它本质上充当了一个优化阻断器或屏障。像指令合并（instcombine）这样的优化传递会知道这个内建函数是“禁区”。

它是规定性的，而非描述性的。我们是在告诉编译器要做什么，而不是向它描述要做什么。

让我们回顾之前看到的那个存在漏洞的Go代码示例，但这次使用了我们的内建函数。它进入LLVM流水线，经过那些优化传递后，生成了如下的x86汇编代码。

; 使用内建函数后生成的汇编
; ... 无数据依赖分支的恒定时间代码 ...

正如你所见，这里没有发生分支或数据依赖的分支。我们只是生成了与之关联的x86恒定时间代码。

我们内建函数的一大优点是：对于拥有原生条件移动（conditional move）或条件选择（conditional select）指令的架构（如x86_64, arm64），我们直接使用这些指令。对于没有此类指令的架构（如arm32），我们则生成位操作，仍然提供恒定时间保证。

实现架构与流水线 ⚙️

我们可以看一下我们已支持架构的流水线，例如x86_64、arm64和arm32。

  1. 代码进入存储（Storage）。
  2. 经过Clang前端处理。
  3. IRBuilder创建与我们的内建函数关联的LLVM内部表示（IR）内建函数。
  4. 进入选择DAG（SelectionDAG）阶段。
  5. 在寄存器分配后（Post-RA）或指令选择后（Post-isel）阶段进行扩展——魔法就在这里发生。

我们在所有优化之后进行Post-RA扩展，因此不必担心优化传递折叠或扰乱我们的代码，从而能够为相应架构生成所需的恒定时间代码。我们也支持Thumb/Thumb1和Thumb2模式。

i386架构遵循类似的流水线，但情况有些特殊。我们有两个阶段：一个自定义选择器（custom selector）来帮助缓解一些标志位（EFLAGS）问题以及类型合法化（type legalization）。但我们仍然在Post-RA阶段进行扩展，以生成位操作。对于arm32也是同样的道理，它没有条件移动指令，因此我们同样生成位操作来提供恒定时间保证。

通用后备方案 🔄

我们实现的另一个优点是，对于那些我们尚未添加原生支持、或者本身就没有原生条件移动/选择指令的架构，我们仍然能提供恒定时间保证。

这条通用流水线采用了不同的策略：我们不是在Post-RA阶段操作，而是在选择DAG阶段实施恒定时间保证。我们进行DAG链式操作，本质上在指令之间创建了人工依赖关系。一旦它进入LLVM流水线并被降低（lower），LLVM便不会触碰这些依赖关系，因此不会破坏它们。

社区反馈与协作 🤝

我们收到了社区非常积极和有力的反馈。之前有过尝试解决此问题的提案，但如前所述，没有一个被上游采纳。我们得到了来自密码学家和LLVM开发者的宝贵意见。

我们还与苏黎世联邦理工学院的研究小组合作，在他们的测试套件中使用我们的内建函数。结果显示，在所有架构上都取得了非常好的成果，同时也表明密码学家现在有了比使用各种古怪技巧更简单的方法来绕过优化问题。

未来展望 🚀

除了C++支持，我们正在研究将这些内建函数添加到Rust语言中。同样，由于Swift使用LLVM，它也可以使用我们的内建函数。鉴于我们已经支持WebAssembly，基于浏览器的密码学也能从中受益。

GCC和Cranelift编译器也可以考虑跟进，引入类似的恒定时间保证机制。

我们还在探索为某些算术操作（例如除法和乘法）提供此类内建函数，并为整个表达式提供恒定时间保证。

总结

本节课中我们一起学习了：

  1. 问题根源：编译器激进的优化会破坏密码学库的完整性，引入时序侧信道漏洞。
  2. 解决方案：__builtin_constant_time内建函数提供了编译器层面的保护，是一种“编译器优先”的安全概念。
  3. 核心优势：它在多种架构上工作，提供可靠的恒定时间保证，得到了社区的积极反馈。
  4. 重要意义：这是在缓解分支类时序侧信道攻击的安全领域迈出的重要一步。

011：支持 LLVM-IR 中的原始数据拷贝的字节类型

概述

在本节课中，我们将要学习 LLVM 中间表示（IR）中的一个新概念——字节类型。我们将探讨当前 LLVM IR 在实现内存拷贝（如 memcpy）时遇到的问题，并了解字节类型如何通过更准确地表示原始内存值来解决这些问题，从而构建一个更正确的编译器。

当前 memcpy 实现的挑战

上一节我们介绍了课程的主题。本节中我们来看看当前 LLVM IR 在实现内存拷贝时面临的根本问题。

考虑 C 语言中的 memcpy 实现，它使用 unsigned char 类型按字节进行拷贝。C 标准保证，任何值都可以被拷贝到一个 unsigned char 数组中，并且我们可以访问和操作该值的对象表示。

在右侧的 LLVM IR 中，通过 unsigned char 类型的加载和存储操作被表示为通过 i8 整数类型的加载和存储。这导致 IR 实现无效，因为 LLVM 中的整数类型不具备 unsigned char 类型的相同属性。

当前 memcpy 实现存在两个主要问题：

  1. 拷贝填充位的问题。
  2. 跨指针拷贝时保留指针来源信息的问题。

接下来，我们将深入探讨每个问题。

问题一：拷贝填充位

上一节我们提到了两个核心问题。本节中我们先来详细看看第一个关于填充位的问题。

考虑一个由 short 和 int 组成的结构体。编译器可能会在两个字段之间插入两个字节的填充以满足对齐要求。

在 LLVM IR 中，这两个填充字节通过一个包含两个 8 位整数值的数组显式表示。目前，LLVM 使用 毒值 来表示填充位，这构成了 memcpy 实现的第一个问题。

目前，整数值在单个值的基础上跟踪信息。这意味着，如果我们尝试将整个结构体作为一个 i64 加载（这在执行加载拓宽优化时经常发生），即使 64 位中只有一位是毒值，整个结构体也会被加载为毒值。这会导致问题，因为我们会将整个结构体作为毒值拷贝到内存中。

问题二：拷贝指针与指针来源

上一节我们讨论了填充位的问题。本节中我们来看看第二个问题，它涉及指针拷贝和指针来源的概念。

首先，简要介绍指针来源的概念。在 LLVM 中，指针并非简单的整数值，它们还包含一些额外的元数据，用于证明优化的合理性。

我们可以将一个指针表示为一个对象标识符和一个偏移值。不同的内存分配会返回指向唯一对象标识符的指针。

以下是描述指针来源的核心概念：

  • 对象标识符：每次内存分配（如 malloc）产生一个唯一的对象 ID。
  • 偏移值：指针算术运算只修改偏移量，保持对象标识符不变。

公式表示：指针 = <对象标识符， 偏移量>

跟踪指针来源的主要优势是简化别名分析，从而启用其他优化。例如，编译器可以假设基于不同对象的指针不会相互别名，即使它们的底层地址可能相同。

然而，当前 memcpy 实现的第二个问题就在于跨指针拷贝时无法保留这种来源信息。

LLVM IR 的当前内存语义

上一节我们了解了指针来源的重要性。本节中我们看看 LLVM IR 当前如何处理内存中的不同类型字节。

目前，LLVM IR 具有类型化内存，意味着我们可以区分两种字节：

  1. 指针字节：除了值之外，还持有指针来源信息。
  2. 整数字节：包含数值。

当前 IR 语义规定，为了避免将指针强制转换为整数并丢弃其来源信息，通过整数类型加载指针会导致结果为毒值。因此，当我们将指针作为整数从内存加载时，结果是毒值。

这意味着，如果我们试图在 memcpy 中拷贝一个指针值，整个指针字节都会被作为毒值加载。

解决方案：引入字节类型

上一节我们看到了当前语义的局限性。本节中我们介绍解决这些问题的方案：为 IR 添加一个新的字节类型。

这个新类型能够按原样表示原始内存值，既可以加载指针也可以加载整数内存，而不会丢弃指针的来源信息。此外，加载带有填充位的值不会污染加载结果，因为字节类型可以在位粒度上维护和表示单独的毒值位。

我们的计划是将 C/C++ 中的原始内存访问类型（包括字符类型和 C++ 的 std::byte）都降低为字节类型。这种新的 lowering 方式使得 memcpy 的 lowering 变得正确，因为我们现在可以存储和加载字节类型，而无需引入隐式强制转换。

字节类型的实现细节

上一节我们介绍了字节类型的概念。本节中我们来看看它的一些具体实现细节。

我们的提案还添加了字节常量，用于初始化全局变量。这些字节常量严格等同于其对应的整数常量。

我们还添加了一个新的 bytecast 指令，允许将字节值转换为其他基本类型。我们允许两种版本的转换：

  1. 默认版本：执行类型擦除，这意味着它通过丢弃来源信息将可能的指针重新解释为整数。
  2. 精确版本（通过指定 exact 标志）：如果字节所持有的值的类型与转换目标类型不匹配，则返回毒值。

代码示例：%val = bytecast b8 %byte_val to i32 （默认，类型擦除）
代码示例：%val = exact bytecast b8 %byte_val to i32 （精确，类型不匹配则返回毒值）

此外，我们还扩展了 trunc 和逻辑右移指令以接受字节操作数，这有利于实现存储转发优化。

字节类型的应用示例

上一节我们了解了字节类型的指令。本节中我们通过一些例子看看它的实际应用。

考虑一个简单的 C 函数，它将 char 值加上常数 1。目前，LLVM 将 char 值降低为 8 位整数，然后进行符号扩展、加法，最后根据 C 标准的要求将结果截断回 8 位整数。

使用字节类型后，char 值被表示为 b8 类型的字节值。然后通过 bytecast 的类型擦除变体将其转换为整数，可能将指针重新解释为其整数表示形式。之后像以前一样执行加法，然后将截断结果转换回字节，因为该函数返回一个 char 值。

字节类型的另一个应用是修复当前由 GVN 等优化执行的一些不安全的值强制转换。通过字节类型，我们可以加载原始内存值而无需引入隐式强制转换，然后使用 bytecast 指令将其转换为整数或指针，从而避免这种不安全的类型擦除。

性能评估与测试结果

上一节我们看了字节类型的应用。本节中我们来看看引入它带来的实际影响。

我们在 SPEC 2017 套件中选取了一组 20 个 C/C++ 应用程序，对解决方案进行了实现和基准测试。

  • 编译时间：未观察到有意义的性能回归（在 -1% 到 1% 之间）。
  • 运行时间性能：未观察到任何重大回归。
  • 峰值内存使用率：变化范围在 -0.5% 到 0.5% 之间。
  • 二进制文件大小：最初发现了一些回归，主要是由循环向量化成本模型的差异引起的。

我们还在 Alive2 中实现了字节类型，并在 LLVM 测试套件上运行。通过将字节类型引入 LLVM IR，我们修复了之前被报告为不安全的测试。这些测试主要由于 memcmp、memcpy 到 load-store 对的不安全 lowering，以及主要由 GVN 执行的不安全值强制转换引起。

总结

在本节课中，我们一起学习了 LLVM IR 中引入的新字节类型。我们探讨了当前 IR 在表示原始内存拷贝（如 memcpy）时遇到的问题，包括处理填充位和保留指针来源信息的困难。

字节类型的引入解决了这些长期存在的问题，使得 IR 能够在不引入隐式强制转换的情况下表示原始内存值。此外，新类型实现了各种内置函数（如 memcpy、memmove、memcmp）的原生实现，并用于修复现有的不安全优化。

我们的实现大约有 2600 行代码，约占 LLVM C++ 代码库的 0.05%。重要的是，字节类型没有导致任何重大的性能回归，并且是迈向更正确编译器的重要一步。

012：关于AArch64上BOLT的教程及其与PGO的竞争或互补关系

概述

在本教程中，我们将学习二进制优化与布局工具（BOLT），重点探讨其在AArch64平台上的应用，并将其与其他基于性能剖析的优化（PGO）技术进行比较。我们将了解代码布局问题、BOLT的工作原理、其与LLVM PGO的异同，以及如何在实际应用中选择合适的工具。

代码布局问题与解决方案

上一节我们介绍了教程的概述，本节中我们来看看代码布局面临的核心问题及其解决方案。

问题：糟糕的空间局部性

假设我们有一个要在处理器上执行的二进制文件。执行开始时，我们从二进制文件加载一些指令到指令缓存中。当遇到分支时，我们跳转到新位置，必须从新位置加载指令到缓存中。如果频繁发生分支跳转，就会导致缓存不断刷新。

然而，并非所有代码都均匀执行。有些函数是“热”的，意味着它们运行更频繁或时间更长；有些则是“冷”的，可能根本不运行或运行时间很短。基本块也是如此，只有少数是热的。当我们的热路径在二进制文件中跨越很长的距离时，就产生了糟糕的空间局部性。本质上，热点代码分散在整个二进制文件中。

如何识别问题

识别此问题的方法是使用硬件性能指标。一种方法是采用自上而下的分析方法。当工作负载受前端限制时，这可能是一个优化的好候选。对于BOLT，来自其作者的经验法则是：当每千条指令的L1指令缓存未命中数超过10次时，就是一个候选；超过30次则是更强的信号。

我们也可以使用性能分析工具，例如在AArch64上可以使用topdown工具，或者使用perf。很快我们将能够使用ATP或Arm Performance这个全新的工具。

为什么编译器需要帮助

让我们看看Clang的编译流程。我们从源代码开始，经过前端生成IR，然后在IR上进行各种优化，接着是MIR，生成机器码和目标文件，最后链接成二进制文件。问题在于，在每个阶段，编译器都只能依赖启发式方法。我们需要关于调用图和控制流图的“情报”来确定热点，这在编译时是无法预测的。

当涉及到链接时优化（LTO）时，流程类似，但会在IR上运行一个更重量级、能力更强的第二遍优化。然而，它仍然面临同样的问题，即使用启发式方法，因此无法帮助代码布局。

显然，我们需要重新排列代码。为此，我们需要性能剖析数据。

解决方案：代码布局与性能剖析

上一节我们明确了编译器需要性能数据的帮助，本节中我们来看看具体的解决方案，主要围绕代码布局和性能剖析两个方面。

代码布局优化

我们可以做三件事来优化代码布局：

  1. 函数重排：并非所有函数都是热的，我们可以根据热度重新排序函数。这样做可以消除那些导致CPU流水线中断的昂贵跳转。
  2. 基本块重排：因为只有少数基本块是热的。一旦我们这样做，本质上就是优化了顺序执行的情况。这意味着一旦进入这样一条链式块，我们很可能会一直停留在这条链上，直到完成才需要跳转到另一条链。这是一个针对每个函数的旅行商问题。
  3. 代码分割：因为在不同的轨迹（即那些热块链）之间跳转时，我们必须跨越包含一些冷代码的距离，这可能再次导致CPU中断。因此，最后一件事是将代码分割成一个薄的热区和一个更大的冷区。在冷代码布局方面我们比较幸运，因为通常只有一小部分二进制文件会主导执行。可以想象，一端是热循环，另一端是异常处理块等，它们应该存在，但很少出现在热路径上。

性能剖析方法

PGO或基于性能剖析的优化，是任何使用性能数据来优化代码的技术。本质上，BOLT也是一种PGO技术。两者以及其他PGO都可以带来显著的加速。然而，它们并不那么完善，可能需要额外的步骤来应用，其中一些步骤（例如BOLT的性能剖析）有不同的实现方式。此外，BOLT所做的和其他PGO所做的之间存在一些重叠，这造成了一些混淆。

PGO远不止代码布局。如果在LLVM中搜索“profile summaries”，可以找到许多匹配项，其中只有高亮显示的与BOLT试图做的事情更密切相关。如果继续查看“block frequencies”和“branch probabilities”，并限制在与寄存器分配相关（与布局无关）且非常具体的方面，甚至可以找到更多匹配项。重点是，PGO远不止代码布局。

以下是三种主要的性能剖析方法：

  • 插桩：修改二进制文件以获得准确的计数。可以针对每个函数或每个基本块。这会非常慢，在生产环境中可能是个问题，并且会增加二进制文件大小，在移动设备上可能是个问题。但至少我们能获得最高质量的性能剖析，这意味着有最高的优化潜力。
  • 基于采样：获取近似的计数，并希望得出相同的结论。在AArch64上，可以使用像BRB这样的单元以及perf等工具。这应该更快，但质量可能参差不齐，因此必须检查是否捕获了足够的信息，否则可能会限制优化潜力。
  • 追踪：使用像ETM或PTM这样的单元，会生成海量数据，这意味着开销，但可以使用技巧来限制这些开销。

性能剖析面临的挑战

环境可能施加一些限制，例如在生产环境中插桩可能是禁止的，或者移动设备上的二进制文件大小可能是个问题。硬件也可能有限制，例如BRB目前在AArch64机器上很难找到。

还需要在性能剖析质量和开销之间进行权衡。插桩质量最好，但开销最高。对于基于采样的方法，我们可以调整这种权衡，例如增加采样频率意味着更高的开销和可能更高的性能剖析质量。我们必须确保对性能剖析的时间开销和优化效果都满意。

当然，工作负载可能会变化。如果今天获取的性能剖析数据，可能下个月就不够好了，或者代码可能发生变化，这可能意味着必须重复性能剖析步骤。

最后一个挑战，也是与BOLT和LLVM PGO比较相关的，是关联和维护问题。考虑这个流程：我们用Clang编译一个二进制文件，然后用perf获取一些性能数据。这些性能数据是通过对二进制文件采样捕获的，因此它们处于二进制级别。当我们开始使用PGO时，我们从源代码开始编译，因此必须经历从二进制级别到源代码IR的关联过程。这个过程可能导致一些性能剖析质量损失，从而可能限制优化潜力。然后，当我们开始优化并经历所有优化过程时，必须维护性能剖析质量。之前的演讲提到了调试信息的一些问题，当我们进行这些转换时，转换越深入，某些转换扭曲或稍微歪曲性能数据的可能性就越高，这可能意味着我们可能限制优化潜力。之前的演讲也提到，我们是人，可能会犯错。例如，代码交换了后继块，但性能数据本身没有交换。结果，此后的任何优化过程都会做出错误的假设。这很棘手。

BOLT详解

上一节我们探讨了性能剖析的通用方法和挑战，本节中我们深入了解一下BOLT这个具体的工具。

BOLT代表二进制优化与布局工具。它由Meta创建，已有十多年历史，并在生产环境中使用了多年。它是一个后链接优化器。从这个流程来看，我们只关心二进制文件。实际上，我们不在乎是GCC还是Clang生成的这个二进制文件。

BOLT需要的只是一个二进制文件和一个性能剖析文件。然后，BOLT可以反汇编并理解这个二进制文件，执行关联过程，将性能剖析数据附加到BOLT的IR上。BOLT的好处在于，这个IR处于汇编级别，它使用机器指令。因此可以说BOLT是一种低级别的PGO，因为它使关联和维护过程变得更加简单直接。因此，即使使用基于采样的方法，它也能拥有高质量的性能剖析数据。

BOLT支持布局，即重排函数、基本块和分割代码。完成这些后，我们就可以得到优化后的二进制文件。它支持之前提到的所有性能剖析方法及其变体（除了追踪）。

与其他后链接优化器的比较

以下是BOLT与其他一些后链接优化器的比较：

  • 输入二进制文件：
    • BOLT：不需要重新编译。我们可以直接获取一个二进制文件并重新布局。但是，要进行完整的布局（包括函数重排），我们需要二进制文件中保留位置信息。好消息是这些位置信息只占用少量额外空间，因此将其加入优化流程不是大问题。
    • Code Locality：需要每个函数有独立节区的目标文件。
    • Propeller：需要每个基本块有独立节区的目标文件。
  • 重排能力：
    • Code Locality：只能重排函数。
    • BOLT和Propeller：可以进行完整的重排（函数和基本块）。
  • 编译器支持：
    • BOLT和Code Locality：支持LLVM和GCC。
    • Propeller：需要LLVM工具链，集成到构建系统中可能需要更多努力。
  • 二进制重写：
    • BOLT是一个重写工具，因为它反汇编二进制文件，理解它，然后重新链接。这些额外知识给了BOLT优势，使其可以运行一些额外的优化。但它是在二进制级别运行，源代码IR仍然是某种汇编，所以我们不能用它完全替代LLVM级别的优化。
    • 其他工具不具备此能力。但有些用户可能认为不重写是优势，因为他们信任链接器进行原始布局和链接，而使用BOLT则需要信任BOLT进行重新链接。

LLVM PGO 详解

上一节我们介绍了BOLT，本节中我们来看看LLVM生态系统中的PGO是如何工作的。

Clang本身使用启发式方法，结合PGO则使用性能剖析数据。加上LTO，我们会有一个额外的过程来进行更好的过程间优化。

在性能剖析技术方面，我们有基于采样的PGO（也称为AutoFDO、SPGO等）。在AArch64上，可以使用分支栈（BRB）。BRB是一个分支记录缓冲区，允许我们有效地对调用图中的控制流进行采样。它之所以优化，是因为它只捕获分支，这样我们可以在后处理中推断出前向路径。当然，它也支持插桩，可以在前端级别或IR级别进行。还有上下文敏感的变体。

上下文敏感是什么意思？举个例子：函数A调用函数foo 20次，并且每次都调用函数p。函数B调用函数foo 10次，并且每次都调用函数bar。如果没有上下文敏感，我们只知道foo被调用了30次中的20次。但有了上下文敏感，我们可以假设，例如，A最终在所有调用中都调用了p，这对于某些特定工作负载可能非常强大。

LLVM如何进行代码布局

  • 函数排序：发生在链接器级别。使用PGO时，lld会自动为我们完成，但我们也可以使用顺序文件自己完成，例如Propeller和Code Locality工具就是这样做的。
  • 基本块重排：发生在MIR管道的各个阶段。即使没有性能数据，它也必须运行，因为你必须发出这些块。但是当我们有性能数据时，它会使用这些数据运行旅行商问题的近似算法以做得更好。
  • 代码分割：发生在MIR管道中仍然可行的最后阶段。默认情况下是关闭的，无论我们是否使用PGO。根据我的理解，使用CSPGO加上LTO或SPGO，如果我们不想使用BOLT，我们可以最接近BOLT的效果。BOLT会拒绝重新布局这样的二进制文件，这也意味着如果没有分割，我们可以在LLVM PGO之上应用BOLT。

LLVM PGO 与 BOLT 的比较

在LLVM级别，我们处于更高级别，可以做的事情多得多，因此它当然比BOLT强大得多。但这带来了之前提到的关联和维护问题。特别是基于采样的方法可能受关联问题影响，因为我们在二进制级别捕获性能剖析，然后必须开始用源代码IR进行编译。包括基于插桩在内的任何方法都必须经历维护过程，这可能会损失一些性能剖析质量。我们越深入这些管道，增加这种歪曲或扭曲的可能性就越高。请记住，代码布局发生在管道最后期的阶段。这就是BOLT的好处：它在需要使用性能剖析数据的地方拥有高质量的数据。

因此，当人们问哪个更好时，我认为我们有两个不同级别、解决不同问题的工具：高级PGO和低级PGO。当人们问应该使用哪一个时，答案应该是两者都用。例如，可以先使用IR PGO加上LTO，然后在上面应用BOLT。

其他PGO技术

还有一些其他的PGO技术，其中一些更针对特定工作负载，可能不支持大多数目标平台。

  • CSPGO：上下文敏感采样PGO。它结合了之前看到的分支栈和调用栈来提供上下文敏感性，这是两个特性中的第一个。第二个特性（我认为更通用）是使用伪探针，这会产生最小的运行时开销，本质上为每个基本块附加一些唯一标识符，这极大地帮助了关联过程。因此，通过类似伪插桩的方式，我们可以提高性能剖析的质量。
  • 硬件PGO：是SPGO的扩展，使用更多硬件单元的数据。引入此功能的作者使用分支预测信息来判断一个分支是否真正不可预测。有些信息即使通过插桩也可能难以获得。因此，当一个分支真正不可预测时，他们采取了不同的收敛路径。我最近看到的另一个技术（我认为仍然属于硬件PGO范畴）使用内存性能剖析来优化热路径中的虚拟负载。
  • 数据局部性：我们不仅可以做代码局部性，还可以做数据局部性。例如，我们可以分割数据节区。
  • 时间局部性：不仅仅是空间局部性。例如，对于移动应用程序，这些应用程序可能甚至不受CPU限制。我们关心的是更快地加载应用程序。因此，我们尝试以一种在应用程序运行之前需要最少内存页面的方式布局，这将意味着更少的页面错误和更少的手机启动画面时间。

AArch64 平台演示

上一节我们理论介绍了各种PGO技术，本节中我们通过一个在AArch64平台上的实际演示来看看BOLT的效果。

演示准备

我从一个病理性的案例开始。最初，我们有一个二进制文件，其中热点代码分散在整个二进制文件中，导致CPU流水线中断。这很可能是许多数据中心工作负载的情况，或者是一些更复杂的二进制文件的情况，因为编译器默认无法进行良好的布局。

对于这个演示，我想要一个简单的、编译快速、没有头文件、甚至可以使用插桩、没有依赖关系的程序。所以我选择了冒泡排序，这也是GCC使用的一个例子。但就其本身而言，它对BOLT来说不够好（或不够差）。所以我做了一些修改让它更有趣：我把交换部分包装成一个函数；然后添加了一些带有开关的控制流（这些开关从不运行），所以BOLT必须发现这一点并把它们推到一边；我做了几个副本；然后在它们之间添加了更多函数。所以BOLT必须找出这些黄色的热点部分，把它们拼接在一起，并把所有东西推到某个冷区。这仍然只有大约100KB，不是一个巨大的二进制文件。但请注意，这并不是典型情况，这个二进制文件具有极差的局部性。因此，我只是在这里演示用法，实际工作负载的结果当然会有所不同。

演示步骤

  1. 编译程序：使用你喜欢的编译器编译这个程序。我使用GCC -O3，并保留位置信息，因为我们想要进行完整布局。我们确认一下对象文件中的函数顺序，确实如源代码所示。
  2. 确认问题：在优化之前，我们运行topdown工具确认这是一个有问题的二进制文件。可以看到前端限制达到了55%，这是由于每千条指令有60次L1指令缓存未命中，并且也有相当多的分支预测失误。
  3. 使用插桩优化：
    • 首先对二进制文件进行插桩：bolt --instrument --instrumentation-file=profile.fdata input_binary -o instrumented_binary
    • 运行原始二进制文件和插桩后的二进制文件生成性能数据。插桩后的运行速度慢得多，符合预期。
    • 应用BOLT进行完整布局：bolt --data=profile.fdata --reorder-blocks=ext-tsp --reorder-functions=hfsort --split-functions --split-all-cold input_binary -o bolted_binary
    • 运行优化前后的二进制文件进行对比，优化后的运行时间大约减少了一半。
    • 再次运行topdown确认改进：前端限制从55%下降到30%，L1指令缓存未命中降至每千条指令少于1次，分支预测失误也从大约16次降至10次。
  4. 使用BRB性能剖析：使用perf record -j any_call捕获分支记录。由于BRB不广泛可用，我使用了从另一台机器获取的BRB性能数据。使用perf2bolt将perf数据转换为BOLT可识别的格式，然后应用BOLT。
  5. 使用PEBS性能剖析：使用perf record -e br_inst_retired:near_taken捕获分支指令。同样使用perf2bolt转换，然后应用BOLT。
  6. 使用基本PMU采样：使用perf record -e instructions捕获指令采样。转换后应用BOLT。

性能剖析方法对比

  • PMU：提供单个指令指针。
  • SP：提供一对分支地址，无法像BRB那样展开。
  • BRB：可以展开，理论上（也很可能在实践中）需要更多的开销来捕获与BRB相同质量的性能剖析。因此，BRB是最佳选择。
  • 如果没有BRB，也不介意插桩，可以使用它。也可以尝试ETM/PTM，也许配合一些过滤来管理开销。或者可以使用SP或PMU。理论上，甚至可以增加PMU采样频率以达到接近插桩的优化效果，但这取决于工作负载。

BOLT统计信息与可视化

BOLT有一些统计信息来指导用户。例如，基于某些性能剖析实现的代码密度指标。当我们想用BOLT优化时，我们期望二进制文件中的某些冷区会变热，如果我们没有在该二进制文件中找到一些密集区域，我们假设需要捕获更多性能剖析。

还有一些连续性指标，试图判断性能剖析是否合理。它们基于三个观察：

  1. 控制流上的连续性：一个接收到采样的基本块，必须存在一条从入口点能够传递这些采样的有效路径。
  2. 调用图守恒：例如，调用foo 10次，那么foo的入口点应该将这些采样进一步传播到函数内部。
  3. 控制流守恒：一个非入口、非出口的基本块接收到采样，这些采样应该在该函数内进一步传播。

最后一个工具是Hitmap，它帮助我们将性能剖析可视化，以查看二进制文件中采样的分布情况。我们需要一个二进制文件和一个perf性能剖析文件。在演示中，我们展示了优化前后二进制文件的Hitmap。优化前，热点分散；优化后，所有采样都集中在代码的一个小区域，直观地展示了优化效果。对于一个真实的大型数据库二进制文件的缩放视图，优化前可以看到颜色和采样分散在各处；优化后，许多颜色和采样都紧凑在一个小的热区，而滚动查看时，样本很少，没有颜色，可能是小写字母，并且有更大的空白间隙，这表明我们做得很好。

总结与团队介绍

在本教程中，我们一起学习了BOLT工具及其在AArch64平台上的应用。我们探讨了代码布局问题、BOLT作为后链接优化器的工作原理、其与LLVM PGO的差异和互补关系，以及各种性能剖析方法。关键点在于，BOLT在二进制级别操作，能够避免高级PGO中的关联和维护问题，从而在代码布局优化方面可能获得额外收益。通常的推荐是结合使用LLVM PGO（用于高级优化）和BOLT（用于低级别布局优化）。

我们是Arm的一个团队，致力于BOLT的开发。我是AArch64后端的维护者，我们尝试通过补丁、审查和测试提供帮助。我们非常感谢LLVM上非常活跃的社区。当然，也要感谢Meta创造了这个伟大的工具，这是他们的成果。

BOLT效果显著。可以尝试使用topdown工具查看是否有好的候选对象，然后应用BOLT。如果需要帮助，可以通过Discord、课程页面、LinkedIn联系我们，我们每两周的星期三也有办公时间。今天晚些时候（4:15）还有一个圆桌讨论。

本教程的版本将上传到提供的链接。我必须根据这些幻灯片来撰写它。

问答环节

问：关于使用BOLT的要求。如果我使用Swift（它有LLVM后端和PGO标志），BOLT能开箱即用吗？还是需要一些额外工作？
答：官方支持C/C++二进制文件和PIE。如果二进制文件本身，你可以尝试看看会发生什么，但不会得到官方支持。我知道有些人用它处理Go和Rust二进制文件，这些也不受官方支持。

问：典型的加速效果如何？对于一个符合经验法则、想要使用BOLT的程序，典型的改进是什么样的？
答：这取决于具体情况。我们有加速达到40%的案例，那是一个特定程序（某些函数被内联了数千次）。对于已发布供使用的二进制文件，在某些情况下效果显著。当这些问题得到改善后，加速效果当然会下降。当在PGO之上应用BOLT时，加速也会下降，因为你需要PGO来获得良好的初始布局，需要LLVM进行内联等优化，然后应用BOLT来获得可能由于关联和维护问题而丢失的几个额外百分点。理论上我们不需要BOLT，但实际上我们确实能获得一些额外的百分比。

问：BOLT的二进制重写对调试信息有何影响？是否需要维护或更新？
答：是的，需要维护。实际上，我们有一个专门针对调试信息的容器，但它必须被支持。我个人没有过多关注调试信息。我们最近遇到了一些与DWARF相关的问题，现在有工作在进行以提供支持。我认为是基本支持。

问：对于大型应用程序中跨多个DSO（动态共享对象）的前端性能跳跃，BOLT能否解决？它优化单个DSO，是否有范围解决这种复合问题？
答：我认为如果你在运行程序时捕获单个性能剖析，你可以获得所有DSO的性能数据。然后你必须对每个DSO单独应用BOLT。对于DSO本身在内存中的放置，很可能取决于ASLR。在缓存方面，我认为这不太重要，因为地址的最后部分可能用于映射。但我们可以进行一些改进以避免许多冲突。也许可以改进为在单个命令行中处理多个DSO，而不是手动操作。

问：我尝试了BOLT，但性能没有变化。topdown显示有10%的错误推测和20%的前端限制（在x86上）。如何检查我的二进制文件是否适合BOLT优化？
答：如果在尝试BOLT前后进行topdown分析，可以了解是否发生了变化。BOLT还有其他统计信息（如--dyno-stats），可以显示它尝试做了什么。结合性能数据和实际优化情况，可以告诉你，例如，优化前有多少分支被采用，优化后因为优化了顺序执行而减少了多少。这可能是另一个表明二进制文件是良好候选的迹象。但其中一些是经验法则，可能需要反复试验。建议通过提供的渠道联系并提供具体示例，有人会尝试帮助。

问：对于大型（例如GB级别）的可执行文件，BOLT是否会显著增加编译时间？
答：是的，随着代码规模增大，BOLT处理时间会增加，因为它需要处理更多代码。但如果你尝试对那个二进制文件进行PGO，也需要更多时间。BOLT将遍历代码，而PGO则需要经历所有优化过程。这是特定于代码大小的。

问：性能收集和优化阶段之间，程序差异的容忍度如何？是否需要每次更改程序都重新收集性能剖析？
答：是的。Meta增加了对稳定性能剖析的支持。这意味着如果你有一个二进制文件并收到一些小的更新，它可能能够匹配之前获取的相同性能剖析，并在不重新剖析的情况下编译这个二进制文件。但这取决于二进制文件变化有多大以及匹配效果如何。如果二进制文件变化或工作负载变化，你可能需要重新剖析。在极端情况下，你可能需要为每个输入优化一个二进制文件。一种常见的做法是获取不同的性能剖析，然后合并它们，BOLT可以合并不同的性能剖析文件，然后得到一个应该对所有输入都足够好的性能剖析。理论上，你可以针对不同输入进行不同的优化。

问：BOLT相比Propeller有哪些特定优化的优势？除了代码布局，数据布局优化有好处吗？
答：BOLT的主要用途是代码布局。但根据你的工作负载，可以尝试用BOLT做一些事情。它进行的一些优化可能在某些情况下是额外的收获。例如，最近添加的一个功能也可能在LLVM中实现。我个人没有单独尝试查看每个优化过程的效果。在AArch64上，我们的主要目标是与x86看齐。我们正在尝试收集更多信息，添加更多指南，以便用户知道有哪些选项可用，并可能提供一些统计数据来支持它们。这会很有趣。

013：libc++近期性能工作概述

📋 概述

在本教程中，我们将学习libc++标准库近期在性能方面的工作。内容涵盖：微基准测试基础设施的改进、新的基准测试工作流程、具体的优化案例、性能随时间追踪的工具，以及LNT性能追踪工具的现状与未来。

🧪 1：微基准测试基础设施

libc++长期以来一直拥有基准测试，但过去它们并未集成到测试套件中，需要通过常规CMake手动构建和运行，这带来了不便。

去年，我们将这些基准测试集成到了libc++的Lit测试套件中。在过去一年里，基准测试的数量大幅增加，目前约有11,000个微基准测试。每个基准测试都针对特定的API、特定的使用模式和特定的数据大小进行测试。例如，向一个包含8,000个元素的map中插入元素就是一个独立的基准测试。

这些基准测试使用Google Benchmark库编写。编写方式与编写普通的Google Benchmark类似。例如，下面是一个对32位整数进行std::hash基准测试的示例：

static void BM_HashInt32(benchmark::State& state) {
    std::vector<int32_t> data = generate_random_data(state.range(0));
    for (auto _ : state) {
        for (auto val : data) {
            benchmark::DoNotOptimize(std::hash<int32_t>{}(val));
        }
    }
}
BENCHMARK(BM_HashInt32)->Range(8, 8<<10);

开发者只需在libc++测试目录中创建一个以.bench.cpp为扩展名的文件并编写测试。Lit测试运行器会自动识别这是Google Benchmark测试，并确保正确的依赖项和链接设置。

运行这些基准测试与运行其他libc++测试一样简单。在后台，测试运行器会多次执行代码以确保测量结果的稳定性，并生成LNT兼容的输出文件。LNT是一种性能追踪工具，其输入格式非常简单：<基准测试名称> <度量标准> <值>。

为了方便本地迭代，我们还提供了一个简单的脚本，用于比较两个LNT格式文件，并给出几何平均差异等有用信息。

🔄 2：改进的基准测试工作流程

上一节我们介绍了如何本地运行基准测试。本节中，我们来看看如何让这个过程对第三方贡献者更加友好，特别是如何实现部分自动化。

我们最近添加了一个GitHub Actions工作流，它可以通过PR评论触发。贡献者只需在PR评论中输入特定命令（例如 libcxx-benchmark <基准测试名称>），GitHub Actions就会在后台启动，在PR的基准分支和主分支之间运行A/B比较，并将结果以评论形式回复。

目前这个功能仍处于实验阶段，我们暂时没有专用的硬件来运行它，因此测量结果存在一定噪音。我们复用了用于测试的预提交CI环境，该环境并未针对基准测试进行优化。但基础设施已经就位，未来可以在此基础上进行改进。

⚡ 3：近期优化案例

现在，让我们结合上述工具，看看libc++近期的一些具体优化案例。自LLVM 20以来，我们落地了超过25项有记录的优化，性能提升是当前22发布周期的重点。

我们通过分析内外部工作负载，调查了调用频繁且存在改进空间的API。其中，关联容器（如std::map, std::set）脱颖而出，成为重点优化对象。

优化 std::map 的复制操作

std::map 本质上是一棵红黑树。过去，复制一个map的方式比较低效：创建一个空的目标map，然后逐个插入源map中的所有元素。每次插入都需要从根节点开始重新查找插入位置。

然而，源map本身已经是一棵结构良好的红黑树。优化思路是直接复制这棵树的结构，而无需进行任何键值比较。这项优化为map的复制构造和复制赋值操作带来了 10% 到 90% 的性能提升。

优化 std::map 的查找操作

查找操作是map的核心。标准map使用“严格弱序”比较器（通常是<）。在查找时，对于树中的每个节点，算法需要执行两次比较：

  1. 判断查找的键是否 小于 当前节点的键。
  2. 如果不是，则判断当前节点的键是否 小于 查找的键（等价于判断查找的键是否 大于 当前节点的键）。

对于某些重要类型（如std::string），它们提供了“三路比较”操作（如<=>），可以在一次操作中确定小于、等于或大于的关系。

我们的优化是：当map的键类型是std::string且使用默认的std::less比较器时，在查找算法内部使用三路比较来替代两次<比较。这直接将每次节点访问的比较次数减半。

这项优化为std::map<std::string, ...>的find、count、contains等基本操作带来了 30% 到 50% 的性能提升。而对于std::map<int, ...>等类型，由于没有应用此优化，性能保持不变。

📊 4：追踪libc++性能随时间的变化

性能优化有时会意外引入性能回退。因此，持续追踪性能随时间的变化至关重要。它可以帮助我们及时检测和修复回归，验证优化的实际效果，甚至发现历史上引入但未被察觉的回归问题。

我们最近引入了一个脚本，允许在任意指定的Git提交上构建和测试libc++。该脚本不会污染当前的Git工作区，而是在临时目录中进行操作。结合新的基准测试运行方式，我们可以回溯历史，针对旧的libc++代码运行当前的基准测试套件，从而收集历史性能数据。

例如，我们可以获取2025年1月以来的所有提交，并针对每个提交运行map的基准测试。脚本会为每个提交生成一个LNT格式的输出文件。我们还有另一个本地工具，可以将这些数据可视化，生成一个交互式的HTML图表，清晰地展示每个基准测试随时间的性能变化，并能关联到具体的提交信息。

🛠️ 5：LNT工具的状态与未来

上一节提到的本地图表工具，其功能与LNT（LLVM Nightly Test）工具的核心目标一致。LNT是一个更完善的性能追踪系统，采用客户端-服务器架构。服务器集中存储、可视化数据并自动进行回归检测；客户端工具则负责在任意架构和配置上运行测试并提交数据。

LNT的现状是：之前运行的实例（lnt.llvm.org）已停止服务，代码库也有一段时间未积极维护。但目前复兴工作正在进行中，我们修复了大量单元测试，预计在几周内重新部署一个可用的LNT实例。

未来的计划包括：

  1. 重新部署 lnt.llvm.org。
  2. 将本地比较工具的功能合并回LNT，方便社区使用。
  3. 进行UI改进，例如在图表上显示提交信息。
  4. 改进回归检测算法。
  5. 利用LNT追踪尽可能多的历史基准测试数据。
  6. 追踪多种配置，如libc++的硬化模式、稳定/不稳定ABI，以及在多种架构上的表现。

一旦基础设施就位，我们将系统性地调查历史上的性能回归，这些都是唾手可得的性能提升机会。

🎯 总结

本节课中，我们一起学习了libc++在性能方面的一系列工作：

  1. 改进的微基准测试基础设施：将大量基准测试集成到Lit套件中，方便运行和比较。
  2. 更友好的工作流程：通过GitHub Actions实现PR级别的自动化基准测试比较。
  3. 具体优化案例：包括优化std::map的复制和查找操作，后者通过利用三路比较获得了显著提升。
  4. 性能追踪工具：介绍了用于历史性能数据收集和可视化的新脚本与工具。
  5. LNT的复兴：讨论了这款强大性能追踪工具的当前状态、即将到来的部署以及未来规划。

最终目标是使libc++不仅比现在更快，而且通过正确的工具和持续的努力，力争成为最高效、性能最强的可持续C++标准库实现。

014：高通基于MLIR的技术一瞥

概述

在本节课中，我们将学习高通公司如何利用MLIR框架，为其Hexagon NPU构建一个全新的机器学习编译器。我们将了解其整体工作流程、核心优化策略，以及如何应对大模型编译的挑战。

章节 1：编译器架构与目标

上一节概述了课程内容，本节中我们来看看高通MLIR编译器的整体架构与设计目标。

我们正在构建一个基于MLIR的编译器。尽管高通内部有一些遗留的编译器，但我们正在基于从这些遗留编译器中学到的经验，以及MLIR开源社区的先进成果，构建一个全新的、基于MLIR的编译器。

我们主要关注三个领域：

  1. 构建编译器核心：这意味着从我们的入口方言（linelang）开始，逐步将其降低到LLVM IR，然后使用后续的工具链。
  2. 前端支持：我们目前专注于PyTorch和Triton。PyTorch用于完整模型编译，Triton用于内核编写。
  3. 硬件目标：目标是Hexagon NPU。

到目前为止，我们已经取得了非常令人鼓舞的成果。我们拥有一个功能完整的端到端工作流，可以展示LLM（如GPT、Llama）的完整模型编译与执行。我们也成功编写并映射了一些从简单到非平凡的Triton内核到我们的硬件上。我们认为，达到未优化代码性能的70%到80%是一个良好的开端。

章节 2：硬件特性与工作流程

上一节我们介绍了编译器的架构，本节中我们来看看目标硬件Hexagon NPU的特性以及对应的编译工作流程。

Hexagon NPU是一种VLIW（超长指令字）多线程硬件。它包含标量单元、向量单元和矩阵单元。它拥有专用内存、向量寄存器和暂存器。编译器的复杂性在于需要妥善处理所有这些硬件特性，确保正确编排所有计算和数据移动，并充分利用硬件提供的优势。

以下是我们为PyTorch和Triton设计的工作流程：

我们专注的核心代码是MLIR Hexagon编译器，它接收linelang输入并生成LLVM IR。我们所有的优化都发生在这个代码的Pass管道中。为了实现端到端流程，我们依赖社区工具链。例如，对于PyTorch，我们使用Torch-MLIR转换器将PyTorch模型转换为linelang。对于Triton，我们是开发Triton共享中间层的少数团队之一，该层是Triton到linelang的转换器。然后，我们使用编译器代码将linelang降低到LLVM。

需要指出的一点是，我们的关键方法是构建自己的、针对特定硬件的优化，但同时重用并利用社区中发生的所有技术进步。我们也利用高通内部多年开发的手写Triton内核库。如果存在手写内核，我们就利用该库；如果不存在，我们就自动生成代码。因此，这是一种自动代码生成和重用现有手写库的混合模式。

章节 3：核心优化策略

上一节我们了解了工作流程，本节中我们来看看编译器实现的核心优化策略。

我们构建的许多优化都是针对特定目标的，旨在利用我们的硬件特性，并根据硬件约束来编排计算和数据移动。我们的方法是在上游可用组件的基础上进行构建。

我们进行的优化包括融合（Fusion）和分块（Tiling）。分块出于多种原因：你希望将计算分布到并行单元上；你需要确保计算的数据占用适合你的暂存器容量；分块也有助于向量化等。

数据移动对性能至关重要。你需要确保使用正确数量的缓冲区，将它们放置在正确的位置，并以能够适时预取的方式进行数据移动，确保计算过程中成本得到分摊。因此，融合、分块、向量化、数据移动等编译器中常见的优化一直是我们的重点。

此外，我们拥有用于向量矩阵计算的特殊硬件单元，它们需要特殊的数据布局。我们也利用上一讲演讲者提到的linelang打包（pack）和解包（unpack）操作来进行数据分块，确保数据被转换成我们硬件所需的布局。利用上游方言，但将其转化为我们硬件能够使用的方式，这就是我们的方法。

以下是一个简明的示例说明：

  • 在左侧，你可以看到代码经过一定程度的融合，以最大化带入暂存器的数据。
  • 融合完成后，进行分块。分块出于多种目的：分配并行性、使数据适应暂存器、确保最内层循环能够向量化。
  • 同时，为创建的每个分块，将数据移入缓冲区，并在适当时机将数据移出缓冲区，放置DMA操作以实现数据预取。

所有这些都需要妥善编排和执行，以获得优化代码。在这个过程中，我们重用了许多上游方言，但也创建了自己的目标特定方言。

另一个来自Triton的示例：我们尝试获取他人编写或我们自己编写的Triton代码，这些代码对目标硬件是通用的。然后使用我们的编译器，将任何通用的Triton代码映射到硬件上，并在此过程中执行我刚才谈到的所有优化。

此外，正如之前提到的，存在许多已经编写多年的库。这个示例具体展示了，如果你有一个手动优化的库，Pass管道可以确保映射到这个手动优化的库，而不是自动生成代码。因此，拥有这种功能可以兼得自动代码生成和使用现有库的优势。

章节 4：面向大模型的优化扩展

上一节我们讨论了核心优化，本节中我们来看看为增强编译管道可扩展性，以应对当前及未来大模型的一项优化。

我们最近必须开发这项优化的起点是：对于大型语言模型，我们希望能够将模块拆分成多个部分，由操作系统独立加载。也就是说，我们不是只生成一个大的二进制文件，而是希望生成一堆二进制文件并分别加载它们。

为了实现这一点，考虑到超过99%的内存占用来自于常量参数（通常是权重、偏置等），如果我们要拆分模块，就需要将常量与代码分开降低。

我们已在MLIR层面将其实现为一个Pass。这个Pass的工作方式如下：它从一个包含代码和常量的原始MLIR模型开始，然后将其拆分成 N+1 个模块。第一个模块（模块0）只包含代码，而模块1到N只包含数据（即常量参数）。

这个MLIR Pass有点特殊，因为它不是一对一的转换。它接收一个输入模块，但产生一堆输出模块。它主要负责两件事：

  1. 动态创建多个MLIR模块。
  2. 修改原始代码模块，主要是将 I.constant 操作替换为引用其他模块中定义的常量的操作。这些常量模块将被独立编译。

接下来需要做一些工作来适当地编译和链接所有这些模块。可以想象，第一个代码模块需要链接到所有其他常量模块，因为它使用了在其他模块中定义的常量。但对于常量模块来说则更容易，它们不需要链接任何依赖项。

最终，我们能够将一个模块拆分成一堆共享对象文件，并在Hexagon DSP使用的QuRT操作系统上独立加载它们。

章节 5：总结与展望

在本节课中，我们一起学习了高通公司基于MLIR为Hexagon NPU构建机器学习编译器的技术概览。

我们正在高通构建一个面向ML模型的、基于MLIR的编译器，目标硬件是Hexagon NPU。正如Matt所说，我们既利用了开源社区的进步，也借鉴了高通内部遗留编译器的经验。

我们能够使用同一个编译器来降低PyTorch模型和Triton内核，并且已经能够为多种模型（如GPT系列、Llama系列等）生成代码。我们目前正在提高所生成代码的效率。

目前，我们仅支持一个NPU，但我们计划在未来支持多个NPU。更重要的是，我们将把所有成果开源并贡献到上游，预计在今年12月发布。欢迎大家前来与我们交流。

总结

本节课中我们一起学习了高通基于MLIR的Hexagon NPU编译器架构，其融合PyTorch与Triton的工作流程，针对硬件特性的核心优化策略（如融合、分块、数据移动编排），以及为支持大模型而设计的模块拆分Pass。该编译器旨在兼得自动代码生成与手写优化库的优势，并计划开源以贡献社区。

015：用安全的Rust驯服GPU编程

概述

在本教程中，我们将探讨如何利用Rust语言的安全特性来简化并增强GPU编程。我们将从Rust安全与不安全代码的核心差异出发，分析GPU编程面临的主要挑战，并介绍一个基于LLVM的解决方案——offload。我们将详细讲解其工作原理、当前实现状态以及如何利用Rust的所有权系统来实现安全的数据并行访问。

章节 1：Rust中的安全与不安全代码

Rust语言通过借用检查器（borrow checker）强制执行严格的内存安全规则，这能防止大量错误代码，但也限制了某些操作。另一方面，这种限制为我们向LLVM中间表示（LLVM IR）添加属性和传递信息带来了信心。

上一节我们提到了Rust的安全哲学，本节中我们来看看安全与不安全代码在LLVM IR层面的具体差异。

安全Rust函数

一个简单的安全Rust函数定义如下，其中包含两个指向f64类型的引用（&f64）：

fn safe_function(x: &f64, y: &f64) -> f64 {
    *x + *y
}

从该函数生成的LLVM IR中，可以看到参数x和y都带有noalias属性，以及其他一些对齐（align）和noundef信息。这表明编译器可以假设这两个指针不会指向重叠的内存区域。

不安全Rust函数

不安全Rust通过unsafe关键字启用额外功能，如解引用裸指针（raw pointer）。这对于底层编程或与其他语言交互非常有用，但编译器无法提供安全保障，需要开发者自行确保代码正确性。

unsafe fn unsafe_function(x: *const f64, y: *mut f64) -> f64 {
    *x + *y
}

从不安全函数生成的LLVM IR中，noalias属性消失了。这意味着使用不安全Rust可能会错过一些性能优化机会。与C/C++中的restrict关键字不同，Rust没有直接等效的关键字。

核心概念：安全Rust的引用在LLVM IR中默认带有noalias属性，而不安全Rust的裸指针则没有。这是影响性能的关键差异。

此外，即使是在安全Rust中，也存在一些特殊的“魔法”类型（如UnsafeCell），它们允许在拥有不可变引用的情况下修改底层数据，在这种情况下也可能不会生成noalias信息。

章节 2：GPU编程的核心挑战

从上一节对Rust内存安全模型的讨论中，我们已经可以窥见GPU编程面临的主要挑战。本节我们将通过一个具体的向量加法示例来分析这个问题。

以下是CUDA编程指南中一个经典向量加法的Rust直译版本：

// 理想化的安全Rust GPU内核（目前无法编译）
fn vector_add(a: &[f32], b: &[f32], c: &mut [f32]) {
    let idx = thread_index(); // 获取当前线程索引
    if idx < c.len() {
        c[idx] = a[idx] + b[idx];
    }
}

这里的问题在于Rust严格的别名规则：你可以同时拥有多个对同一数据的不可变引用，但只要存在一个可变引用，就不能有任何其他引用（可变或不可变）。这是为了在LLVM层面保证noalias。

显然，我们希望内核在多个线程上并行运行，并且需要将计算结果写入某个可变的状态（c）。如果尝试用普通的Rust编译器编译这段代码，或者通过类似Rayon的库进行并行化，代码将无法通过编译，因为借用检查器会阻止这种并行可变访问。

那么，我们如何才能安全地实现GPU编程呢？

章节 3：Rust GPU编程基础设施现状

在深入探讨解决方案之前，我们先了解一下Rust目前为GPU编程提供了哪些基础设施。

以下是当前Rust生态中与GPU相关的主要目标和支持状态：

  • NVIDIA GPU：通过nvptx64目标支持，目前为第2级（Tier 2）支持，意味着可以编译，但功能有限。标准库（core）的一部分已被移植，提供了一些内部函数，如获取线程ID和共享内存操作。
  • AMD GPU：通过amdgpu目标支持，目前为第3级（Tier 3）支持，状态更不完善。之前甚至存在编译问题，例如尝试为不支持的目标生成f128浮点数。
  • rustc代码生成器：Rust编译器内部有一个GPU内核ABI，它会根据编译目标（NVIDIA或AMD）选择对应的ABI，并进行基本的健全性检查，确保Rust代码适合该目标。

总的来说，当前的基础设施处于可用但功能有限的阶段。

章节 4：目标与现有方案

我们为Rust GPU编程设定的目标是：安全、便捷、支持原生Rust类型和多后端。具体来说：

  1. 默认安全便捷：利用Rust的所有权信息，实现自动内存传输。
  2. 提供不安全逃生舱：允许有经验的开发者为了性能或功能进行底层控制，例如手动管理内存传输。
  3. 支持原生Rust类型和函数：旨在支持绝大多数Rust函数和类型在GPU端运行。
  4. 支持多厂商后端：不局限于单一硬件平台。

基于这些目标，我们评估了现有方案：

以下是现有Rust GPU编程方案的对比：

  • RustaCUDA：仅支持NVIDIA，并且由于之前提到的别名问题，它迫使用户到处编写不安全代码。
  • cuberuby等DSL方案：大多在Rust之上实现了一个领域特定语言，绕过了编译器，因此不支持原生Rust类型。
  • offload（我们的方案）：基于LLVM，能够支持多厂商（目前支持NVIDIA和AMD）。它底层依赖于经过OpenMP和C++/Fortran测试的特性，有望在合理时间内稳定。由于基于LLVM，它可以与其他有趣的特性（如基于Enzyme和LLVM插件的自动微分）协同工作。此外，通过集成GPU libc项目，我们甚至可以在GPU上进行基本的I/O操作，这对调试非常有帮助。

我们选择了offload作为实现路径。

章节 5：offload方案的工作原理

现在，让我们深入探讨offload方案的具体实现。我们将从用户接口开始，逐步了解其编译流程。

用户接口

目前，我们首先实现一个普通的CPU函数。例如，一个对包含256个f32值的数组进行操作的函数：

fn cpu_function(data: &mut [f32; 256]) {
    for i in 0..256 {
        data[i] = data[i] * 2.0;
    }
}

为了在GPU上运行它，我们提供了一个核心内部函数offload。它首先注册函数名，以便在后续编译步骤中识别哪些函数需要为GPU目标编译，然后转发参数。

// 使用 offload 内部函数
#![feature(offload_intrinsics)]
use std::intrinsics::offload;

fn main() {
    let mut data = [0.0f32; 256];
    unsafe {
        offload::kernel_launch(cpu_function, &mut data);
    }
}

我们计划扩展这个接口，目前诸如线程块大小等参数是硬编码的，未来应该能方便地随参数传递。

编译流程

当前的编译流程分为两步，并且还依赖一些实验性功能。
以下是offload的编译步骤：

  1. 编译主机代码：使用 cargo build --target <host-target> 或直接调用 rustc。
  2. 编译设备代码：直接调用底层编译器 rustc，为其提供源码，并启用实验性的 #![feature(offload_intrinsics)] 和 -Z unstable-options。我们指定GPU目标（例如 --target nvptx64-nvidia-cuda 或 --target amdgpu）。
  3. 链接与打包：目前仍需手动使用 clang-offload-packager 和 clang-offload-wrapper 进行链接和打包，我们正在努力自动化这一过程。

为了稳定此功能，我们需要消除对实验性功能的依赖，例如稳定 build-std 功能，或者将AMD和NVIDIA目标提升到更高级别的支持等级。

章节 6：实现安全的数据并行访问

回到最核心的安全性问题。我们之前看到，通过多个线程并行访问可变引用 &mut 是未定义行为。offload 如何解决这个问题呢？

我们有两种主要策略来修复这个特定的用例。

策略一：通过裸指针和手动索引

由于裸指针没有别名保证，我们可以安全地在多个线程间共享。每个线程获取自己的线程索引，然后基于该索引对裸指针进行偏移计算，最后从偏移后的裸指针创建一个安全的引用。

// 简化示例：每个线程处理一个元素
fn vector_add_unsafe(a: *const f32, b: *const f32, c: *mut f32, len: usize) {
    let idx = thread_index();
    if idx < len {
        // 对裸指针进行偏移计算，不通过它进行读写
        let c_ptr = c.add(idx);
        // 从偏移后的指针创建唯一引用
        let c_ref = unsafe { &mut *c_ptr };
        *c_ref = unsafe { *a.add(idx) } + unsafe { *b.add(idx) };
    }
}

关键点：即使可变引用 c_ref 和原始裸指针 c 指向的数据有重叠，但只要我们不通过原始裸指针 c 进行读取或写入，这段代码就是安全的（sound）。我们只是用 c 来做地址计算。

策略二：创建非重叠切片（Slice）

如果每个线程负责处理输出数组中的一个连续子区间，我们可以为每个线程创建一个独立的切片（&mut [T]），这些切片在编译时或运行时确保互不重叠。

// 概念性示例：将数组分块给不同线程
fn process_chunk(data: &mut [f32], chunk_id: usize, total_chunks: usize) {
    let chunk_size = data.len() / total_chunks;
    let start = chunk_id * chunk_size;
    let end = start + chunk_size;
    let my_slice = &mut data[start..end];
    // 安全地处理 my_slice
}

显然，这种索引逻辑可以被自动生成。Rust编译器拥有所需的所有信息：数组长度、切片大小以及线程数量。

章节 7：处理复杂索引与不安全抽象

然而，我们无法期望编译器永远能自动推断出正确的索引模式。用户可能需要更复杂的索引逻辑，而不仅仅是简单的逐元素或连续块处理。

根据对 rust-perf 基准测试的分析，大约三分之二的测试用例可以通过简单的标量或批量索引逻辑覆盖。对于剩下的部分，我们需要更灵活的机制。

这就是不安全抽象（unsafe abstraction）的用武之地。我们提供一个不安全的接口，让用户自己定义如何将输入数据划分为不重叠的部分，同时确保其划分逻辑在语义上是正确的（例如，确保索引是单射的，即不同的输入索引映射到不同的输出位置）。

以下是 rust-perf 基准测试中遇到的复杂索引模式示例：

  • 标量索引：大多数简单情况。
  • 基于查找表的索引：使用一个索引 i 从数组 A 中查找值 j，再用 j 作为索引访问数组 B。这需要确保 A 中的值是单射的，否则会导致数据竞争。
  • 多重嵌套循环：目前尚未覆盖。

对于这些复杂情况，用户可以通过实现特定的 unsafe trait 来提供自定义的、经验证的数据划分逻辑，并将其封装在安全的接口之后共享给社区。

章节 8：未来优化方向

目前，每次 offload 调用都会产生三次数据移动（到设备、内核启动、回主机）。我们计划进行多项优化以提升性能。

以下是计划中的主要优化方向：

  • 按需拷贝：利用“默认不可变”的特性，很多情况下数据无需拷贝回主机。
  • 设备端直接分配：对于常见场景（如预先分配大量GPU内存），支持直接在设备上分配变量，避免在热路径中分配。
  • 内核间数据驻留：如果数据在主机端未被使用，则让其保留在设备上，避免不必要的来回拷贝。这需要分析主机代码中的数据流。
  • 共享内存支持：目前仅通过实验性方式在AMD端暴露，需要与社区合作使其稳定可用，并探索如何提供尽可能安全的接口。
  • 内核融合：将多个连续的内核调用融合，减少启动开销和数据传输。

此外，为了支持现有的生态系统（许多线性代数库在底层使用不安全代码），我们正在探索一种方案：通过检查这些库实现的 Copy/Clone trait 的LLVM IR，并将其中的内存拷贝操作替换为设备与主机间的拷贝。这在大多数情况下应该是安全的。

章节 9：总结与问答摘要

本节课中我们一起学习了如何利用安全的Rust来进行GPU编程。

内容总结

  1. 安全与不安全Rust：安全Rust通过借用检查器和LLVM IR中的noalias属性提供内存安全保证和优化潜力；不安全Rust则提供更多控制但牺牲了部分编译器保障。
  2. GPU编程挑战：Rust严格的别名规则与GPU数据并行编程中需要多线程写入同一数组的需求存在根本冲突。
  3. offload方案：我们介绍了一个基于LLVM的offload方案，它支持多后端、原生Rust类型，并旨在提供安全便捷的默认体验。
  4. 安全并行策略：通过使用裸指针进行地址计算，然后创建唯一的引用或切片，可以实现安全的数据并行访问。编译器可以自动处理简单情况，复杂情况则通过不安全抽象由用户提供验证过的划分逻辑。
  5. 未来工作：重点在于性能优化（如减少数据移动、支持共享内存、内核融合）和生态集成（支持现有不安全库）。

问答摘要

  • 关于共享内存的安全性：目前认为很难提供完全安全的共享内存接口，因为无法防止用户引入数据竞争。可能最终会将其标记为 unsafe。技术上，可以通过OpenMP类似的接口（如额外的 dyn_ptr 参数）来提供支持。
  • 关于动态索引的支持：对于运行时决定的、非单射的索引模式，无法保证安全，必须通过不安全接口。在某些情况下可以加入运行时检查。
  • 关于内核间数据驻留：基本思路是跳过中间的数据回传操作，只在所有offload调用开始和结束时进行整体数据传输。这需要分析数据在主机代码中是否被使用。
  • 关于线程数据所有权的推理：核心是通过裸指针计算偏移，但不通过该裸指针进行读写。然后从偏移后的地址创建引用，这个引用是唯一的。只要遵守“不通过原始共享裸指针进行访问”的规则，即使引用与裸指针指向的区域重叠，也是安全的。

通过本教程，我们希望为你展示了在Rust中实现安全、高效的GPU编程不仅是可能的，而且能充分利用Rust语言本身的优势来构建更可靠的系统。

016：Clang中的生命周期安全性分析

概述

在本教程中，我们将学习Clang中引入的生命周期安全性分析。该分析旨在编译时发现C++程序中的“释放后使用”等时序内存安全问题。我们将了解其核心概念、工作原理以及如何帮助开发者编写更安全的代码。

问题背景：时序内存安全

上一节我们介绍了本教程的主题，本节中我们来看看该分析旨在解决的核心问题。

时序内存安全问题是指：访问已被释放或回收的内存是一种未定义行为。这通常表现为悬垂指针，并可能导致程序崩溃、数据损坏甚至安全漏洞。

其影响不容小觑。例如，在Chromium项目中，超过三分之一的关键安全漏洞是由“释放后使用”这类经典的时序安全问题导致的。

让我们快速看一个“释放后使用”问题的例子：

  1. 程序在某个时刻创建了一块内存（可能在栈、堆或由栈对象管理的堆上）。
  2. 在该对象存活期间，程序获取了指向它的指针（别名）并将其存储在某个变量中。
  3. 在之后的某个时间点，该内存被释放（例如，对象离开作用域）。
  4. 在此之后，任何通过之前持有的指针访问该内存的行为都属于“释放后使用”，是未定义行为。

分析概述

上一节我们了解了问题的严重性，本节中我们来看看解决方案的总体思路。

生命周期安全性分析试图在编译时发现这些时序内存安全问题。这是一种基于别名的分析，属于“左移”方法，旨在在开发早期识别问题。

该分析的目标是：

  • 提出一个直观的生命周期模型。
  • 通过渐进式类型等机制，实现增量式的时序内存安全保障。

核心分析维度

上一节我们介绍了分析的目标，本节中我们将深入探讨程序员如何推理“释放后使用”问题，这引出了分析的三个关键维度。

考虑之前的例子，程序员通常会关注以下几点：

  1. 对象在何时被销毁？（验证操作）
  2. 有哪些指针指向了刚被销毁的对象？（别名分析）
  3. 这个指针（悬垂指针）的值后续是否被使用？（活性分析）

以下是每个维度的详细说明：

1. 无效化操作

我们首先关心的是对象何时被销毁。更广泛地说，我们关心的是验证哪些操作会使内存失效。

以下是无效化操作的例子：

  • 普通的delete操作。
  • 管理堆内存的栈对象在其作用域结束时调用的析构函数。
  • 在代码中手动调用的析构函数（通常很糟糕）。
  • 隐藏在抽象背后的操作，例如函数调用：
    • 向vector执行push_back操作会使之前持有的迭代器失效。
    • clear、insert等容器操作同样会导致失效。
    • 这不是容器操作独有的特性，而是任何C++对象固有的属性。

2. 别名分析

其次，我们关心别名。我们想知道一个存储位置有哪些别名。

在示例中，p = &x;这行代码负责创建指向x所支持存储的别名，并将其存入p。因此，我们说p是x的别名。

别名分析需要考虑控制流和任意赋值。例如：

  • 当我们将指向x的指针q赋值给新指针p时，p和q都成为x的别名。
  • 即使引入任意复杂的控制流（如循环、函数调用、条件语句），我们仍然需要回答“x有哪些别名？”这个问题。

3. 活性分析

最后一个维度是：该指针持有的值是否在后续被实际使用？这可以通过活性分析来回答。

活性本质上是用于判断变量持有的值在程序后续是否会被读取的预计算。它也是流敏感的。

例如：

  • 如果有两个赋值创建了别名，第二个赋值p = &y;会覆盖p之前的值（即指向x的引用）。我们说它“杀死”了旧值，旧值不再存活。
  • 在后续有使用p的情况下，p的新值（指向y）是存活的，而旧值（指向x）则不是。

生命周期模型与术语

上一节我们明确了分析的三个维度，本节中我们将定义一个生命周期模型和相关术语来回答这些问题。

该模型灵感来源于Rust的Polonius借用检查模型，但根据C++的需求进行了调整。它旨在拥有类似Rust的生命周期语义，并可能在未来支持类似Rust的语法。

以下是核心概念：

1. 借用

借用表示从内存位置借用的行为。它由“借用创建于何处”和“借用了什么内存”来定义。

例如，当你将整数的地址赋值给指针时：p = &x;，我们就在右侧创建了一个借用L1，它借用了路径x（x可以是任何存储位置，如栈变量、结构体字段等）。

2. 借用过期

借用会过期。当存储位置失效时，指向它的所有借用都会过期。这代表了之前我们关心的“无效化”操作。

例如：

  • p = &x; 创建了指向x的借用L1。
  • q = &x; 创建了另一个借用L2（即使指向同一存储，但借用地点不同，故为不同借用）。
  • 当x离开作用域被析构时，L1和L2同时过期。

3. 起源

起源是与指针类类型关联的符号化标识符。例如，int* p具有一个起源属性（如origin01），它代表了指针的别名部分。

具体来说，一个起源是一个实体可以持有的借用集合。可以将其视为指针可能引用的所有数据来源。

例如：

  • 初始时，int* p有一个起源O1，不包含任何借用。
  • 当处理p = &x;时，创建的借用L1成为起源O1的一部分。

起源可以包含流敏感信息，并遵循子类型规则进行流动。例如：

  • int* p, *q; 有两个不同的起源O1和O2。
  • 当执行q = p;时，右侧起源O1中的借用L1会流入左侧起源O2。此时，O1和O2都包含借用L1。

4. 起源的活性

如果p先借用自x，然后p又借用自y，新的赋值会覆盖旧的借用。后续的读取使得最新值存活，而旧的借用被“杀死”，不再存活。

定义生命周期违规

上一节我们介绍了模型的基础构件，本节中我们将它们组合起来，定义什么是生命周期违规。

在程序点P，如果有一个借用L在点B过期，并且一个起源包含该借用L，同时该起源在点P是存活的，那么我们就说发生了生命周期违规。

更具体地说，生命周期策略是：一个存活的起源绝不能包含一个已过期的借用。

在示例中：

  • 在作用域结束时，借用L1过期。
  • 起源O1包含了借用L1。
  • 起源O1是存活的。
    因此，这构成了一次生命周期违规。

实现与诊断

该分析的简易版本已在LLVM上游实现，可通过某些标志启用。它能够提供类似Rust的三点诊断：

  1. 借用点：识别有问题的借用是在何处创建的。
  2. 无效化点：识别该借用是在何处失效的。
  3. 使用点：指出是哪个使用使得该起源保持存活。

处理抽象（函数调用）

上一节我们讨论了单函数内的分析，本节中我们来看看如何处理函数调用带来的复杂性。

函数调用可能引入新的别名，也可能使其通过参数接收的存储失效。

一种方法是过程间分析，但这扩展性不佳，且函数定义并非在编译期始终可用（可能直到链接时）。

因此，我们需要一种组合性好的分析，能够独立分析每个函数。这可以通过让函数表达其生命周期契约、别名契约和无效化契约来实现。Clang注解和API契约等语言扩展已用于此目的。

1. 函数与别名

函数可以通过参数接收引用并返回其别名。目前，可以使用clang::lifetimebound等一系列注解来表达有限的别名契约，但尚无法表达所有复杂的别名关系。

2. 函数与无效化

函数调用可能引入无效化操作，常见的例子是容器的push_back、clear、insert等。目前尚无标准方法告知函数会无效化哪些内容，但未来可能引入类似clang::invalidates的注解。

注意：此分析功能仍在开发中，请关注LLVM 20.6及以后的更新。

分析的范围与限制

在结束之前，需要澄清此分析不是什么：

  • 它不是C++的严格时序内存安全解决方案，仍然可能写出“释放后使用”的bug。
  • 它是一种机会主义的缺陷发现方法，旨在实现增量式的时序内存安全。
  • 它不是C++提案N的实现。

总结

在本教程中，我们一起学习了Clang中的生命周期安全性分析。我们从时序内存安全问题出发，探讨了该分析旨在解决的三个核心维度：无效化操作、别名分析和活性分析。接着，我们介绍了基于借用、起源和活性概念的生命周期模型，并定义了生命周期违规的条件。我们还了解了该分析如何通过诊断指出问题，以及如何处理函数调用带来的挑战。最后，我们明确了该分析的定位和限制。这项分析是迈向更安全C++代码的积极一步，目前仍在积极开发中。

相关资源：

  • RFC: H6291
  • 双周会议（周三），已加入LLVM日历
  • GitHub项目与标签
  • Discord频道（用于讨论生命周期安全）

017：LLDB MCP 概述与核心概念 🧠

在本教程中，我们将学习 LLDB MCP（模型上下文协议）服务器。我们将了解它如何作为桥梁，连接人工智能应用与 LLDB 调试器，从而实现通过 AI 来控制和查询调试会话。

开发者工具与 AI 集成的两种范式

近年来，在将开发者工具与人工智能集成时，我观察到两种主要范式。

第一种是 AI 作为功能。在这种范式下，人工智能增强了工具现有的能力，但并未从根本上改变工作流程。你熟悉和喜爱的工具变得更好，有时也可能稍差一些。一个例子是尝试通过查看调用栈来自动解释程序崩溃。

第二种范式随着像 Cline、Codeium 和 Copilot 这样的编码智能体而出现，即 AI 作为协调器。在这种范式下，AI 自主行动，并代表你调用各种工具。对于 LLDB 而言，一个例子是要求一个智能体“单步执行这个循环，直到它发生一些特殊情况”。

起初，我对 AI 持怀疑态度。LLDB 的核心原则是“调试器永不撒谎”，这与 AI 的概率性本质难以调和。但我逐渐认识到，当人们与这些编码智能体交互时，他们有不同的期望，并且意识到这些智能体可能会出错。因此，对于 LLDB 来说，这被证明是一条更有前景的道路。

什么是 MCP？🔗

MCP 代表模型上下文协议。它是一个开放标准，用于连接 AI 应用与外部系统（如开发者工具）。其目标是提供服务器与客户端之间的统一接口，这样它们就不必为彼此实现自定义的支持。

如果你觉得上一句话很熟悉，可能是因为你在 LLVM 开发者大会上听过我关于 LLDB DAP 的演讲，我当时说过非常类似的话。DAP（调试适配器协议）为调试做了同样的事情，LSP（语言服务器协议）为语法高亮和代码补全做了同样的事情。现在，MCP 为集成 AI 做了同样的事情。

MCP 服务器的三大核心概念 🧩

一个 MCP 服务器包含三个核心概念：资源、工具和提示。

  • 资源 是静态数据源，为 AI 应用提供上下文信息。我喜欢把它们看作是模型可以读取的“文件”。
  • 工具 是可执行的函数，AI 可以调用它们来执行操作。我喜欢把它们看作是模型可以调用的“API”。
  • 提示 是可复用的模板，有助于构建交互。它们就像是模型可以用来完成更复杂任务的“工作流”。

LLDB MCP 服务器实现了资源和工具。接下来，让我们详细看看这两者。

LLDB MCP 的工具实现：command 工具 🛠️

对于 LLDB，我们暴露了一个名为 command 的工具。其定义包含名称、简短描述和一个 JSON 模式（为简洁起见，幻灯片中已省略）。command 工具接受任何 LLDB 命令，运行它，并返回输出。

这被证明是 LLDB 的一个极佳匹配，因为 LLDB 本质上是文本化的，并且大多数大型语言模型已经在我们的命令上进行了训练。我曾尝试创建多个工具（例如，一个用于单步执行，一个用于设置断点），但通过实验发现，这种单一 command 工具的方式效果最好。

LLDB MCP 的资源实现：调试器与目标 🎯

LLDB 有调试器和目标的概念。在命令行中使用时，你通常有一个调试器，其中包含一个目标。但当将其作为库使用时，你可能拥有多个调试器，每个调试器都有自己的目标。

因此，我们必须决定如何通过 MCP 暴露这个层级结构。最初我考虑使用工具，但最终确定为资源，因为这些实体本质上是应用状态。这意味着客户端可以指定对哪个资源进行操作（例如，操作哪个调试器或哪个目标），或者也可以将决定权交给模型，模型将使用启发式方法找出最相关的一个。

连接与传输模式 🔌

与 LSP 和 DAP 类似，标准输入/输出是默认的传输模式。客户端将服务器作为子进程启动，然后通过标准输入和标准输出进行通信。对于大多数客户端，你只需要告诉它在哪里可以找到 LLDB MCP 二进制文件，一切就设置好了。

当你想让 LLM 完全管理调试会话时，这种方法效果很好。但我们也想支持另一种情况：你已经在一个调试器中交互，然后想使用 AI 来询问它问题。例如，你在 Visual Studio Code 中使用 LLDB DAP，并想询问关于当前暂停位置的问题。

为了支持这种情况，你可以使用新的 protocol server 命令从 LLDB 内部启动一个 MCP 服务器。与标准输入/输出不同，它会创建一个套接字并监听连接。我们还为 LLDB MCP 添加了发现机制，以便自动找到这类调试会话。目前这仅限于单个会话，但在未来，我们希望让 LLDB MCP 能够复用多个 LLDB 实例，并将它们呈现为单个 MCP 服务器。

总结 📝

在本教程中，我们一起学习了 LLDB MCP 服务器。我们了解了 AI 集成工具的两种范式，认识了 MCP 协议及其三大核心概念（资源、工具、提示）。我们深入探讨了 LLDB MCP 如何通过一个通用的 command 工具暴露 LLDB 功能，以及如何通过资源来管理调试器和目标层级。最后，我们介绍了两种连接模式：适用于 AI 全权管理的标准输入/输出模式，以及适用于嵌入现有调试会话的套接字模式。

请尝试使用 LLDB MCP。如果你有任何问题或建议，欢迎与我交流或在线联系。

018：用于C++的未定义行为附录

大家好，欢迎。我的名字是 Shafi Yamore。我是英特尔的一名 Clang 前端开发人员，负责支持 C++。我也是 C++ 标准委员会的成员，主要在核心工作组工作。你们可能通过我在社交媒体上每周发布的 C++ 小测验认识我，也可能通过我的博客认识我，我每年大约发布四次博客，内容主要涉及 C++ 边缘案例，有时也涉及 LLVM 和通用软件开发。

你们也可能通过我的 C++ 冷笑话认识我。今天我有一个笑话要讲，如果你们喜欢，可以尝试在社交媒体上关注我。8个字节走进一家酒吧，酒保问：“我能为你们做点什么？”它们说：“给我们来个双份的。”😊。好的。谢谢，谢谢大家的热烈掌声。

我们今天在这里讨论 C++ 的未定义行为和 IFNDR 附录。UB 代表未定义行为。未定义行为是糟糕的代码。基本上，标准规定，对于这种代码，你无法预期其行为。优化器会假设你不会编写这样的代码，并据此进行优化，这时糟糕的事情就会发生。你知道，安全问题会出现，调试会话会试图找出服务器为何崩溃，总之对每个人来说都是坏事。IFNDR 代表“病式，无需诊断”。在标准中，“病式”部分意味着需要给出诊断信息。而这里的“NDR”部分，指的是我们希望它是病式的代码，但由于技术原因，实现可能无法可靠地诊断它。因此，你可能会遇到很多这样的情况：代码可能编译，可能不编译；可能在某些实现上编译，在其他实现上不编译；可能在不同实现上有不同的行为。所以它不像 UB 那样具有无限制的运行时不良行为。但总的来说，它是不希望出现的行为。很多 C++ 开发者对此理解不深。所以它只是一个陷阱。我经常用 UB 来指代两者，这只是作为一种简写。

好的，让我们从高层次谈谈这个问题。C++ 标准中有很多未定义行为和 IFNDR。下一个问题是它们通常很难识别。看看这个列表，仅针对未定义行为。标准文档有 2600 页。我们可以谈论未定义行为，说程序的行为是未定义的，行为是未定义的，一堆其他词是未定义的，我们可以继续说下去。你可以坐在那里说，我们为什么不直接查找“未定义”这个词。不幸的是，我们在许多其他与未定义行为无关的上下文中也使用“未定义”这个词。所以，如果你想找到所有的 UB 和 IFNDR，你必须通读整个文档，一个一个地找，这是一个乏味、漫长的过程，我做过，需要很长时间。

下一个问题是它们通常很难理解。标准文档不是教程，它是一个规范。语言通常非常密集。通常没有示例，有时有，有时没有，不一致。所以如果你不理解，不明白未定义行为或 IFNDR，你就只能靠自己了。最后，问题是有时 UB 是隐含的。标准规范为符合规范的代码设定了指导方针和规则，有时当你违反这些规则时，它就是 UB。标准通常会明确指出，但有时不会。所以你会说，好吧，至少我们可以通读标准，统一措辞，让一切变得清晰简单，但这可能是一件好事，但它并没有解决问题，因为标准仍然不是教程，你仍然没有理解 UB 的方法。它没有用通俗的语言表述，我们也没有一个合适的地方放置示例。老实说，我们只希望有一个地方能列出所有的 UB 和 IFNDR，你可以一次性看到所有内容，而不必搜索这 2600 多页的文档。

你可能会问，什么是附录？ISO/IEC 指令第 2 部分实际上给出了一个很好的描述。C++ 标准委员会是一个 ISO 委员会，通常称为 WG21。它有很多关于如何运行标准机构以及如何编写标准的文档。它告诉我们，附录用于向文档主体提供额外信息，并且出于几个原因而制定。通常当信息过于冗长，放在文档主体中不合适时；为了区分特殊类型的信息；或者为了呈现与文档中特定应用相关的信息。这听起来像是我们想要的，我们想要一个地方来列出所有的 UB 和 IFNDR，用通俗的解释和示例。需要澄清的是，在标准术语中，我们经常听到规范性和资料性。规范性是规则，当你看到标准中的这些词时，你必须遵守这些规则。资料性通常是我们标准中的注释，它们不是规则，只是为了指导进一步的理解。

那么这里的高层目标是什么？第一个目标是为 UB 创建一个资料性附录。为 IFNDR 创建一个资料性附录。为每个条目提供通俗的解释和代码示例。我们基本上想要两个独立的文档，因为它们是不同的东西，IFNDR 和 UB，尽管它们具有相同的特性，即不直观且是开发者容易掉入的陷阱，但它们是不同的东西，我们不想把它们混在一起。另一个目标是我们有很多不同的目标受众。我们有普通的 C++ 开发者，可能是大学生，也可能是从事这项工作很长时间的专业人士。他们想了解什么是 UB，想避免这些陷阱，想学习。然后我们有培训师，培训师可能是公司内部培训师，公司付费请人来培训员工，也可能是大学教授，C++ 会议或其他会议的演讲者。他们希望有材料能够解释 UB 和 IFNDR，有好的例子，能够教人们避免什么。工具开发者，也就是我们，我们编写工具，编写消毒剂，编写静态分析工具，编写诊断程序。所以我们可以把这个看作一个目标列表，我们是否捕捉到了所有我们能捕捉到的东西，我们是否更好地诊断了它们，我们是否作为消毒剂捕捉到了它们。对于工具开发者，我们希望有这样一个好的列表，你可以直接去查看，然后说，嘿，听着，我们可以在哪些方面做得更好。显然，对于安全分析师也是如此。UB 常常导致安全漏洞，这在科技新闻中经常出现。所以我们希望安全分析师能够有一个列表来查看，能够理解未定义行为，能够向用户解释，能够教用户，能够编写更好的工具来捕捉和诊断未定义行为。

另一个目标是我们想要诊断所有显式的 UB。如果我们遗漏了任何东西，那将无济于事。所以它必须是一个完整的列表。我们故意不记录隐式的 UB，因为这可能容易出错，因为它是隐式的，我们可能不理解这些措辞的意图，所以我们可能会把那个 UB 定义得太窄或太宽。因此，我们希望有一个单独的过程，我们需要去审视并说，好吧，我们有隐式的 UB，我们需要让它变得显式，这样我们就可以在附录中记录它。最后，我们希望确保 WG21 编写提案的过程包括从附录中添加和删除 UB 和 IFNDR。所以每当你写一个提案来改变标准中的内容时，你会有措辞放入标准中，以指定你在做什么。我们希望确保这明确成为编写提案过程的一部分。

我们可能会问，为什么这对 C++ 很重要？我认为毫不奇怪，我们的用户想要一种更安全的语言。随着我继续解释，我们可以看到这将如何最终帮助我们获得一种更安全的语言。UB 是不知情的开发者的陷阱。如果你不知道要避免什么，不知道 UB 和 IFNDR 是什么，你就无法避免它们。总的来说，我们发现大多数开发者不太理解 UB。他们可能听说过，可能知道一些例子，但他们并不真正了解它的全部范围。说实话，专家通常也不容易识别 UB。我无法告诉你我在聊天、小组或会议中有多少次，当你坐下来看这段代码时说，那是 UB 吗？我不知道，也许你知道。所以即使是专家，我们也需要这份文档来帮助我们更好地识别和学习 UB。这也很重要，因为一套通用的示例使 UB 更具体，有助于我们理解这一点。你可以描述各种未定义行为，但有时看到代码才能真正让人明白。所以当你看到代码时，有人会说，哦，等等，我以前见过那段代码，也许我写过那样的代码。因此，能够看到示例很重要。最后，如果我们不记录它，我们就不知道它是否在变得更好。理想情况下，从长远来看，我们希望看到标准中的未定义行为减少。但如果我们不跟踪它，不记录它，我们就无法真正判断。😊。

另一个要问的好问题是，为什么这对客户端很重要？毫不奇怪，我们的用户想要一种更安全的语言。我们的用户经常会沟通，写错误报告，进行讨论，问我们问题。你知道，优化做错了什么，或者他们不明白为什么这样优化。这常常归结为 UB。因此，有一个地方列出所有的 UB 并很好地描述它们，将使我们能够更有效、更好地与用户沟通。只需说，嘿，听着，这是 UB，这就是它，这是你需要避免的。作为工具编写者，它给了我们一个目标列表，可以编写更好的工具，编写更好的诊断程序，编写扩展，希望能更容易地避免未定义行为。我们可以更快地实施解决方案。我们已经有很多安全扩展，这里有一些例子，F D，enostr aliasing，Liy plus plus hardening。我要指出，我们，Aaron，Eric Vd，K1 和我，我们有一个关于 Clang 中强化模式的 RFC。我们希望这能成为一个伞形领域，我们可以帮助为 Clang 开发更安全的代码选项。如果你还没有看到这个 RFC，如果你还没有评论，我建议你看看，并添加评论，或者只是让我们更好地理解它。😊。

嗯。我想谈谈这个过程。也许如果你没有参与标准委员会，你不知道标准委员会是如何做出改变的。一般来说，你想改变 C++ 标准，你必须写一个提案。我的第一次尝试是 P1705。你要写一个提案，然后一般来说，标准委员会由许多小组委员会组成。我的目标主要是演进和核心，演进处理语言的演进变化，你想添加功能，想删除功能，想增强东西。演进是你瞄准的地方，所以附录被认为是标准文档的演进变化。核心工作组处理标准中的措辞，所以当你实现东西时，你必须改变标准，核心工作组确保措辞的变化实际上实现了变化的意图，并且其工作方式和规范是正确的。通常你的提案在第一次接触委员会时无法存活，通常你可能需要改进和添加，更清楚地指定东西。所以这是我在这个领域的第二个提案，P3075。最终你被接受了。然后我们到了这一步，即迭代附录措辞。我们有一个针对当前 C++ 草案标准的 PR，我相信其中包含了 UB 和 IFNDR 的完整列表。所以我们正处于只是进行文字润色的阶段。然后一旦完成，我们说，听着，我们把它弄好了，然后我们得到措辞的最终接受，然后我们可以将其合并到草案标准中，然后至少在过程的那个点上我们就完成了。

我想以行动号召来结束。附录需要更多的编辑。我相信我们处于一个非常好的状态，我们有一个完整的 UB 和 IFNDR 列表，但我们需要很多文字润色，也许更好的例子，等等。所以我们需要更多的帮助来实现这一点。我们需要人们帮助我们识别标准中的隐式 UB。我们知道它在那里，我还没有真正坐下来尝试查看并找出它们是什么，但我们需要人们，如果你不在委员会中，就提交缺陷报告，如果你已经在 W21 中，那么你可以在核心邮件列表中开始对话。但我们需要开始这个过程，这样我们就可以把这些隐式的 UB 变成显式的 UB，并希望消除所有隐式的 UB。最后，我想说，我们需要成为我们想要看到的改变。作为实现者和工具设计者，我们可以实现新的诊断程序，帮助指定 U 的行为。你知道，有 UB。我们坐在那里说，是的，好吧，这是预期的，你所做的是未定义行为，但有时行为可能相当疯狂。所以我们也许可以做得更好，尝试限制优化带来的疯狂程度，并希望通过扩展或其他方式总体上消除 UB。谢谢。如果你们有任何问题，我乐于回答。非常感谢这个精彩的演讲。我完全支持这个。未定义行为是一个广阔的领域，有些人会说你不应该有未定义行为，但他们不知道如何测试它，甚至没有一个清晰的规范。这就是为什么我认为这非常好。我有一个问题。你提到你通读了标准，然后找到了那里列出的所有内容，你是如何处理的？这是方法吗？还是有其他想法？所以你必须手动通读标准，找到所有的 UB，我漏掉了一些。Tom 和其他人添加了我漏掉的，因为你知道，当我开始的时候，然后到我整理好的时候，你知道，有些被指定了，等等。所以它有 2600 页，措辞在整个标准中并不统一。所以你必须，有些词可能被分隔开，比如在下一行。是的，完全正确。所以这是一个艰苦的过程。但一旦我们有了附录，我们就不必再这样做了。所以这是一次性的，应该就是这样。好的，谢谢。我以为可能也有一些专家有一些工具，检查分析工具，说，好吧，我们已经有一些列表了。我们有消毒剂，我们有静态分析。但我的意思是，我不认为我们有任何东西能捕捉到所有形式，也许有些你无法轻易捕捉到。但我认为我们可以做得更好，我认为存在我们可以改进的漏洞。好的，谢谢。

有些未定义行为有时被用作编译器的扩展点。是的，这个附录是否区分了“我们根本无法定义这个”的未定义行为和“我们故意不定义这个，因为我们想要”的未定义行为？不，不，它只是用通俗的语言解释它，并给你例子。它不解释意图，不解释用途。我们可能有技术规范，在其中更详细地讨论，讨论可以使用什么工具来检查等等。我还没有开始这个过程，我不认为其他人已经开始这个过程，但我认为这是一个重要的过程。这只是，好吧，我们需要一个地方可以找到所有东西，我们可以教育人们，是的，然后现在当你写提案时，我们必须添加和删除它们。然后我们必须让它成为过程的明确部分。我认为当你坐下来看一个提案时说，你在添加 UB。哦，好吧，我没有意识到。你为什么必须添加 UB？我希望这能有助于减少 UB，或者得到一些更多的反对意见。是的，很好。谢谢，是的。谢谢你的演讲。当你描述那些不知道未定义行为是什么以及它可能出现在哪里的 C++ 开发者时，我首先受到了伤害。是的，所以谢谢。这是一个新颖的努力，尽管还有很长的路要走。你有没有想过 C++26 契约将如何使这个过程更加机械化，甚至找出给定 API 中什么是未定义行为？我没有参与契约，所以我在英特尔的另一位团队成员 Tom Hoorman 实际上一直非常关注契约。我没有关注它。所以我认为有些人对此肯定有非常强烈的意见。契约界的人已经考虑过这个问题，我还没有和他们详细讨论过这个问题。所以我不确定，我的意思是，是的，我认为在这方面有强烈的想法，但我不知道契约界的人在这个领域说了什么。谢谢。是的，不用担心。谢谢。嗨，我有一个问题，基本上是这样的，在最后，我们有一个行动号召，你知道，想法是我们需要更多的人来仔细研究文档，以找出什么是 UB，以便改进附录。你是否看到了一个先有鸡还是先有蛋的问题，即为了创建附录，你必须理解什么是 UB，但为了理解什么是 UB，你必须有一个列出 UB 的附录？那么让新人跟上来的过程会是什么？我的行动号召针对的是不同的人。编辑，如果你想要编辑，你不必是专家，你可以只是，比如说，总体上非常擅长编辑。但是，是的，那些被要求寻找隐式 UB 的人，更可能是专家，已经是委员会成员的人，或者只是非常了解 C++ 的人。所以一个不熟悉标准文档的普通 C++ 开发者，可能不会是他们来做这件事。但我确实想确保，你知道，这里会有很多 C++ 专家，委员会里有一群人，也许他们有想法。也许有人坐在那里说，哦，是的，实际上我考虑过这个，看过它们，我只是有一段时间没想了，他们只是给核心邮件列表发了一封邮件，开始对话，他们甚至不需要做太多。他们只是写邮件，然后核心会接手，因为他们确实关心这个。是的。我的意思是，有时这也是有意的。例如，参数求值的顺序，可以是从左到右，也可以是从右到左。有些是未指定行为。是的，是的，这是实现定义的。所以我认为这些术语也应该，未指定或实现定义。是的，那些是分开的，那些是分开的东西。所以我们实际上，我们确实有实现定义的。我不确定在 C 中，但它们是分开的概念。所以实现定义的行为不是 UB。所以那是一个独立的概念，还有未指定行为。但问题是，实现定义的行为可能会让人措手不及。但 UB 具有这种无限制的影响，而实现和未指定行为没有。所以对于可移植性，它是危险的。但部分是的，当然，当然，肯定是的。是的，是的。好的，那么谢谢。最后一个问题。嗨，我有一个问题。你有隐式 UB 的例子吗？我没有。我没有。我没有深入研究，我没有一个好的隐式 UB 的例子，是的。那么让我们感谢 Shafik 的精彩演讲。

019：MLIR 的正规形式 🧩

概述

在本节课中，我们将要学习 MLIR 中“正规形式”的概念。我们将探讨为什么需要一个比“规范化”更灵活、更明确的机制来描述中间表示的各种状态，以及如何通过定义和验证“正规形式”来优化编译器流程、减少不必要的转换并明确传递前提条件。

问题引入：漫长的编译时间

上一节我们介绍了课程背景，本节中我们来看看一个具体问题。

我在 AMD 担任编译器工程师，负责机器学习编译器相关工作。我曾参与一个大型海洋模拟项目，该项目对气候研究等领域很有用。

为了让模拟代码在硬件上高效运行，我们决定使用编译器。我们有一个项目，它使用 Enzyme 和基于 MLIR 的 Poly 等项目来编译一个庞大的海洋模型，并在超级计算机上运行。

当我尝试编译这个大型代码库时，遇到了问题。代码库规模很大，有数百万行。我启动编译后，等待了很长时间。在等待了两个半小时后，我终止了任务，编译仍未完成。

我开始分析性能。MLIR 有一个 -mlir-timing 选项，可以显示每个过程花费的时间。分析结果显示，在长达约四小时的编译过程中，88% 的时间（约三小时）都花在了规范化操作上。

什么是规范化？

为了理解问题根源，我需要弄清楚“规范化”到底是什么。

我查阅了各种资料，包括 AI 助手、维基百科和相关博客文章。它们给出的定义基本一致：规范化是一种将程序转换到“规范形式”的变换。

那么，下一个问题是：MLIR 的“规范形式”是什么？

我查阅了 MLIR 官方文档。文档提到，MLIR 有一个单一的规范化过程，它会迭代应用所有已注册操作的规范化模式。这是一个“尽力而为”的过程，并不保证整个 IR 都能达到某种规范形式。

文档依然没有明确定义这个“规范形式”具体是什么。实际上，MLIR 社区内部对于什么应该被视为“规范”也存在很多争论。

关键在于，MLIR 本身并没有一个固定的、全局的“规范形式”。MLIR 支持多种方言、操作和属性，谈论一个统一的规范形式意义不大。

从“规范化”到“正规形式”

上一节我们看到了“规范形式”概念的模糊性，本节中我们来看看一个更实用的思路。

我们需要认识到，不同的编译器优化过程可能需要 IR 处于不同的、特定的状态。例如，在循环优化中：

  • 有些优化（如公共子表达式消除）希望循环不变式代码被提升到循环外。
  • 而另一些优化（例如为 GPU 等加速器进行 outlining）则希望循环不变式代码保留在循环内。

在 LLVM 中，人们甚至需要专门编写 Pass 来“撤销”循环不变代码外提。

因此，我们需要的不是单一的“规范形式”，而是多种正规形式。一个程序可以处于形式 A、形式 B，或者两者都不是。编译器变换应该能够声明其所需的前提形式，并能将 IR 从一种形式转换到另一种形式。

我的建议很简单：让我们开始为模块标注其所处的正规形式。

例如，我们可以给模块添加一个属性，声明“本模块中所有循环不变式代码均已提升”。MLIR 设计精良，我们可以为这种属性附加一个验证器。这样，我们就获得了一种额外的验证机制，它能保证 IR 满足某些对特定变换有用的不变式，而这些不变式并非 IR 的根本属性。

正规形式的定义与价值

以下是正规形式的核心定义：

正规形式是一组关于 IR 组件的附加不变式集合。

它类似于数据库中的“范式”，但在这里特指编译器 IR 的约束条件。这些约束之所以重要，是因为它们可以为我们对 IR 所做的假设命名。

我们并非首创。LLVM IR 中已有类似概念，例如：

  • 循环简化形式：为运行循环 Pass 而准备的标准形式。
  • 旋转循环形式：便于进行循环不变代码外提等形式。
  • 循环闭 SSA 形式：确保 SSA 值的生命周期限定在循环内，便于内存分析和并行化。

正规形式可以作为变换的前置和后置条件。一个变换可以声明：“我要求输入 IR 处于‘循环不变式已提升’形式”。如果该形式的验证器已通过，变换就无需再为此做任何检查。

实例：Linalg 方言的正规形式

上一节我们了解了正规形式的抽象概念，本节中我们通过一个具体方言来看看它的实际应用。

以 Linalg 方言为例，它用于表示线性代数运算（如矩阵乘法、张量收缩）。

Linalg 运算有多种表现形式，构成了一个正规形式的演进链条：

  1. 命名形式：例如 linalg.matmul，表示一个具体的矩阵乘法操作。
  2. 泛化形式：通过 linalg.generic 操作表示，它像一个通用的多维循环，描述了数据访问和计算模式。
  3. 分块形式：在循环内嵌套 linalg.generic 操作，便于进行分块优化。
  4. 融合形式：对分块后的操作进行融合，结构上仍保持“循环+泛化操作”的形式，因此可以继续应用同类变换。
  5. 循环形式：最终将 linalg.generic 降级为基本的循环和内存操作，以便进一步降低到低级 IR。

在一个实际的 Pass 管道中，为了在这些形式间切换并满足每个 Pass 的前置条件，我们不得不频繁地插入规范化 Pass，形如：generalize -> canonicalize -> tile -> canonicalize -> fuse -> canonicalize ...

如果我们将每个变换（Pass 或操作）都赋予类型，声明其输入和输出的正规形式，那么我们就可以构建一个清晰的变换图：

  • generalize 变换：命名形式 -> 泛化形式
  • tile 变换：泛化形式 -> 分块形式
  • fuse 变换：分块形式 -> 分块形式（同构变换）

这样，我们就不再需要那些防御性的、频繁的规范化步骤了。

构建正规形式注册表

我们可以设想一个正规形式注册表，类似于方言注册表或接口注册表。

这个注册表可以回答以下问题：

  • 查询路径：“我的 IR 处于形式 A，如何到达形式 C？” 注册表可以给出需要应用的变换序列（例如，先 generalize，再 tile）。
  • 诊断状态：“这个 IR 目前满足哪些正规形式？” 一个 IR 可以同时满足多个不相关的正规形式（例如，既满足某个 Linalg 形式，又满足某个循环优化形式）。

这极大地简化了非平凡的 Pass 管道管理。用户只需要声明目标（例如，“我要分块、融合、填充、然后运行在 GPU 上”），系统就可以根据注册表自动推导出需要插入的必要形式转换步骤，生成完整的、正确的 Pass 管道。

正规形式带来的好处

以下是采用正规形式方法的主要优势：

  • 明确性与可验证性：正规形式是附加的、可验证的不变式集合，不再是“尽力而为”的模糊状态。
  • 承认现有实践：我们实际上已经有多个“规范化器”，只是没有给它们正式命名。正规形式将其具体化。
  • 自动化前提/后置条件：使得 Pass 的编写和管道组合更安全，可以避免 Pass 间的无效规范化。
  • 减少争论：社区无需再争论某个模式是否属于“全局规范化”。每个人都可以定义自己的正规形式（例如，“Alex 正规形式”、“Matthias 正规形式”），并在需要时使用。
  • 解耦方言依赖：目前，方言之间会因为规范化模式产生操作而相互依赖。将特定形式转换放入独立的“正规形式转换”中，可以减少这种人为的编译时依赖。

总结

本节课中我们一起学习了 MLIR 中“正规形式”的概念。

我们首先从一个耗时的编译实例出发，指出了单一“规范化”概念的局限性。然后，我们提出了“正规形式”作为一种更灵活的替代方案，它是一组可验证的附加 IR 不变式。我们通过 Linalg 方言的例子展示了正规形式如何清晰地描述编译降级过程，并探讨了通过正规形式注册表来自动化管理 Pass 管道、明确变换前提条件的愿景。

总而言之，正规形式旨在将 MLIR 中隐含的、零散的“最佳状态”实践，转化为显式的、可命名、可验证和可组合的实体，从而提升编译器的可靠性、性能与开发体验。

020：为Xcode预览提供动力

概述

在本节课中，我们将学习Apple如何利用LLVM的JIT（即时编译）技术来驱动Xcode的预览功能。我们将探讨如何将原本为静态链接设计的程序动态加载并运行，以支持Swift UI代码的实时预览，并深入了解为实现这一目标而对LLVM JIT系统所做的各项改进。

Xcode预览功能简介

Xcode是Apple的集成开发环境。Xcode预览是一项功能，当你的应用程序包含Swift UI代码时，你可以使用预览宏来获取UI的实时预览。其核心理念是，在你编辑UI代码时，能快速获得代码更改对实际运行UI的反馈。这个预览是真正运行的程序，而非解释执行。我们运行你的程序，并将UI反射到Xcode的画布上，因此你可以与之交互，就像使用你的应用一样。

从双重构建到JIT方案

最初，我们通过双重构建来实现此功能。除了常规的调试构建，还会有一个独立的预览构建，其中添加了一层间接层以实现快速函数替换。这种方式效果良好，意味着编辑代码时可以快速交换函数以显示新UI，但也意味着我们需要构建两次。

从去年开始，我们改用基于LLVM JIT的方案。在新方案中，我们从常规调试构建中加载目标文件，并使用LLVM JIT链接器API来添加快速函数更新所需的间接层。这消除了第二次构建，但也意味着LLVM JIT必须能够加载和运行动态链接模式下原本需要静态链接和运行的任何应用程序。

面临的挑战

这些程序可能非常庞大，包含数千个文件、数百兆字节的代码和数据，以及数十万个重定位项。在macOS上，程序通常被拆分为多个框架和库，因此UI代码可能并不全在主可执行文件中。这些程序也可能非常复杂，混合语言编程非常普遍，例如Swift、Objective-C、C/C++等。项目还可能包含预编译的静态归档文件，需要能够链接。我们可能需要启用多个架构切片，并处理改变执行环境行为的非标准授权，如macOS强化运行时。此外，还需要支持有趣的链接器选项和手写或机器生成的汇编代码。

整体架构设计

我们使用的JIT架构设置如下：

  • 我们始终利用ORC（LLVM的JIT库）在跨进程模式下运行的能力。这包括一个包含预览JIT的控制进程和一个独立的执行进程，JIT编译的用户应用程序将在其中运行。这可以避免JIT及其依赖库的细节污染执行进程。
  • 我们仅从调试构建的产物目录中加载目标文件，在预览中没有惰性加载。
  • 我们在JIT链接器中安装了一个覆盖插件，允许我们在目标文件流经JIT时修改它们，以添加快速函数更新所需的间接层。
  • 我们使用自定义内存管理器将代码和数据从控制进程传输到执行进程。
  • 在执行进程中，我们使用一种新颖的动态加载器集成技术，使JIT编译的代码表现得像预编译的一样。
  • 为了处理大型应用程序，我们始终利用ORC并发生成符号的能力，即并发链接多个目标文件。

接下来，我们将更详细地探讨这些部分。

覆盖插件：实现快速函数替换

覆盖插件相对独立。我们能够实现它，是因为JIT链接器API在设计上提供了处理流经JIT的机器代码的极大灵活性。特别是，JIT链接器API可以重命名和添加代码与数据，因此添加间接层很容易。

以下是其工作原理：

  1. 首次看到一个函数时，我们将其重命名，确保其名称不与程序中的任何其他符号冲突。
  2. 然后，我们引入一个使用原始名称的存根，并将存根指针指向原始实现。
  3. 所有引用该函数的代码都将绑定到这个存根。
  4. 在运行时，我们可以根据需要重定向这个存根。例如，要替换函数体，我们可以加载第二个定义（同样会重命名以避免冲突），然后只需更新存根指针。

何时进行函数替换是安全的，这个决定超出了预览JIT的范围，我们将其交给客户端处理。客户端有办法知道在程序仍在运行时是否可以动态替换函数，或者是否需要重启进程。

内存管理器：支持代码签名

我们为此项目对内存管理器所做的更改，主要是为JIT编译的代码应用代码签名。我们这样做并非出于安全考虑，而是在IDE环境中工作，其核心理念是允许执行任意代码。我们进行签名的原因是为了不影响启用了macOS强化运行时的应用程序的行为。

强化运行时是一项可选的授权，它声明进程不应加载未签名的代码。默认情况下，它会禁用JIT，意味着你无法加载没有有效代码签名的库。这对于防止部署应用时的代码注入攻击非常有效，但当我们想要“注入”整个应用程序时，就带来了挑战。

我们通过引入一种与强化运行时兼容的自定义签名类型来解决这个问题，而不是禁用它或寻找变通方法。这种签名类型仅适用于可以附加调试器的程序，因此不会削弱强化运行时的安全性。这意味着当我们为JIT预览加载代码时，我们加载的是有效签名的代码，强化运行时对此感到满意。

JIT链接器API足够灵活，可以轻松表示这一点。我们的内存管理器API允许我们拆分链接后的代码和数据。链接后的数据通过共享内存发送，而为附加代码签名，我们必须创建文件，因为在macOS上，签名附加在文件上。因此，每个链接后的代码段都被写入磁盘作为一个文件，签名后，再映射到执行进程的指定地址。

这部分内容非常macOS特定，但我分享这个故事是因为，如果你曾对使用LLVM JIT感兴趣，但又担心在受限环境中JIT API可能无法工作，那么通常可以通过自定义内存管理接口来处理这类问题，从而使JIT正常工作。

动态加载器集成：核心创新

这个项目中最有趣的部分是动态加载器集成。这源于一个看似边缘情况的问题：如果符号foo在JIT编译的代码中，dlsym(foo)应该返回什么？

对于那些不了解的人，dlsym是一个API，允许你通过名称查找C符号。如果你调用dlsym并传入字符串"main"，你期望它返回指向你的main函数的指针。

答案是，dlsym(foo)应该返回foo的地址。如果它是预编译的，这就是它应该返回的结果，我们不希望这种行为改变。问题在于，这个对dlsym的调用可能本身就在预编译的代码中，这超出了JIT的控制范围。在macOS上，dlsym由系统动态加载器实现，因此我们需要系统动态加载器的帮助。

我们解决这个问题的方法是，教导系统动态加载器将JIT编译的代码视为预编译的代码。你可以将其视为一个伪动态库，其行为由API和回调提供，而不是由动态加载器从磁盘上读取已知格式的文件，并且在这种情况下，回调将由ORC JIT编译器实现。

这样做的好处是，像dlopen和dlsym这样的POSIX API得到了原生支持。但更有趣的是，在这个世界中，预编译的代码可以与JIT编译的代码绑定，这意味着你可以选择性地仅JIT编译单个动态库。

为了理解这对预览用例的优势，考虑一个包含多个库的应用程序，我们想要编辑位于依赖树底部的LibUI库中的UI代码。在旧世界中，预编译代码无法看到JIT编译的代码，我们将不得不JIT编译从根库到目标库的每一个库，这不是因为我们真的想JIT编译它们，而是因为这是它们绑定到下层符号的唯一方式。在这个新世界中，JIT编译的代码对预编译代码可见，我们可以只JIT编译我们想要的部分。性能上的机会在于，你JIT编译正在变化的部分，预编译其余部分，从而获得两全其美的效果。我们永远无法在加载不变的预链接代码方面击败系统动态加载器，但我们可以直接重用它的成果，并将JIT的注意力集中在正在变化的代码上。

如果你想知道，当磁盘上可能没有LibMyUI时，如何静态链接LibBar库，这正是LLVM TAPI（文本API）库解决的问题。它们允许你仅使用一个简单的表（本质上代表一个动态库）来进行链接，而无需拥有完整的已编译动态库。

伪动态库的操作与实现

伪动态库的操作可以归结为几个核心步骤：

  1. 创建：对应register API。它接受一组回调、一个地址范围和一个“可加载路径”谓词。你可以将其视为向系统动态加载器引入了一个虚拟文件系统。
  2. 打开：当系统动态加载器遇到一个它可能想要打开的路径时，如果该路径在磁盘上没有文件，它可以询问已注册的伪库：“你在这个路径可加载吗？”如果其中一个回答“是”，动态加载器就可以说：“好的，运行你的初始化器，你被打开了。”
  3. 查找符号：如果动态加载器需要在伪动态库中查找任何符号，它可以调用查找回调，该回调将转发到我们实现中的orc::lookup。
  4. 关闭：运行反初始化器。
  5. 注销：相当于删除，意味着你不能再打开这个伪动态库。

在ORC运行时中实现这些功能变得简单，因为我们在JIT运行时中已经模拟了这些操作。对于这个项目，我们只需要添加粘合代码来对齐接口，修改内部以添加一些缓存以提高性能，并适应DYLD的锁定方案。我们还添加了自动注册和注销功能，这样当你在控制进程中创建JIT编译的动态库时，相应的伪动态库会立即在执行进程中变为可打开状态。

绑定过程示例

让我们看一个预编译代码如何绑定到JIT编译代码的示例。假设我们正在打开一个名为LibFoo的预编译库，它包含一个需要绑定的外部引用符号bar。bar将定义在一个名为LibBar的伪动态库中，但还没有人查找过它，因此它在进程中还没有地址。bar的定义位于构建产物目录中的一个目标文件中，它已注册到预览JIT，但尚未为其完成任何其他工作。

接下来会发生的是：

  1. 系统动态加载器会遇到对bar的引用，并说：“我需要找到这个符号的地址。”
  2. 它将调用伪动态库上的查找回调，该回调由JIT运行时实现。
  3. JIT运行时将跨进程调用预览JIT，说：“我需要这个符号的地址。”
  4. 像在ORC JIT中一样，符号查找将触发具体化。因此，我们将把bar的定义链接到我们的执行进程中。
  5. 此时，JIT知道了它的地址，可以将其返回给ORC运行时，ORC运行时再返回给动态加载器，动态加载器就可以绑定地址。
  6. 现在，我们可以完成预编译代码的打开，并且预编译代码已经绑定到这个JIT编译的定义上。

性能优化与改进

在性能方面，我们首先启用了并发链接。这之所以容易，是因为我们从2018年就开始致力于并发支持。尽管如此，面对我们处理的代码规模，我们还是发现并修复了一些竞态条件，特别是在macOS平台上。

我们还必须为此引入一个自定义调度器。当JIT发现需要做一些工作时，它不会只在当前线程上运行，而是会将工作交给ORC调度器，这是一个你可以自定义的类。在这个项目中，我们引入了一个具有固定数量链接线程，但无限数量“请求处理”线程的调度器。这是因为ORC和ORC运行时之间的许多请求仍然被表述为阻塞操作，固定的线程池可能会被饿死。我们不喜欢这样，认为这是一个bug，并将动态线程池任务调度器视为一种临时方案。但到目前为止，这是一个经过实战检验的方案。

我们还为此项目改进了许多实用函数。例如，LinkGraph的splitBlock操作在链接器中拆分内容。对于重复应用，它原本是O(N²)复杂度，以前从未成为问题，但面对我们处理的代码规模，它突然变得耗时。我们添加了缓存，使其变为O(N log N)，现在不再是问题。LLVM JIT中的许多其他函数也得到了同样的优化处理。

最大的性能变化发生在一个名为“等待图”的数据结构中。这是允许ORC并发的数据结构，它通过跟踪JIT中每个符号当前正在等待哪些符号来确保查找安全。我们能够通过合并LLVM JIT中的两个旧API——addDependencies和emit——来使其变为二分图。现在，你只在发出符号时告诉JIT它们的依赖关系，并且你可以划分该图，说有一组已发出的节点正在等待一组未发出的节点。这种二分图在实践中比旧的任意图更容易处理。

我们还开始对这些节点应用合并。因此，具有相同依赖关系的节点现在被合并为一个更大的节点。这在实践中非常有帮助，因为对于大型程序，许多符号最终会堆积起来，使用一些我们仍在尝试通过JIT推送的通用API。通过合并这些节点，我们极大地缩小了图的规模。

等待图已被提升为独立的类型，它过去是嵌入在JIT动态库中的。好消息是，它现在终于可以进行单元和性能测试了。此外，节点标签过去是引用计数类型，但我们能够证明，在这个图中的任何东西在别处的符号表中都有相应的引用，因此这些引用计数永远不会降到零，它们是冗余的工作。我们能够采用新的非拥有型SymbolStringPtr类型作为图中节点的标签，从而也消除了引用计数。

总的来说，所有这些改进使得我们一些极端情况下的性能提升了100倍到10000倍。

结果与未来展望

就结果而言，目前只有粗略的数据。大约87%的预览在300毫秒内运行，另外约10%在1秒内运行，约1.5%在2秒内运行，剩余的0.7%运行时间超过2秒。我们对这个初步结果感到满意，但必须谨慎解读这些数字，因为存在一些偏差。例如，新项目用户通常会落入300毫秒的区间，这会抬高该区间的比例；而对于确实存在性能问题、预览速度很慢的项目，用户可能会停止使用该功能，这又会人为地缩小最慢类别的比例。我们认为目前的性能是合理的，但知道仍然存在一些极端情况，性能优化工作将继续进行。

处理各种特殊情况

最后，我想谈谈为了使这个项目正常工作而需要处理的各种特殊情况：

  • 命名归档：系统链接器对文件扩展名很宽松，只要路径解析为可链接的内容，它就会链接。为了匹配这一点，我们添加了一个名为orc::loadLinkableFile的便利函数，它可以处理目标文件、归档文件、通用二进制文件等。
  • ld -r合并的目标文件：这些文件可能包含重复的符号名，而常规编译器输出的目标文件则不会。Swift包管理器会这样做，因此许多项目都会间接遇到这种情况。我们移除了局部作用域符号名唯一的假设，现在支持这种情况。
  • ARM64指针认证：这是ARM64上的一项功能，允许你在指针的高位添加一些比特，硬件能够在有人伪造指针时捕获，这是一种控制流完整性措施。我们现在能够在JIT编译的代码中支持这一点，并且有趣的是，我们能够在不引入通用原语的情况下做到这一点。
  • 还有许多其他功能现在也得到了支持，如紧凑展开信息、C++异常、通过符号的子节、弱加载和隐藏链接、静态归档链接的-all_load和-ObjC修饰符，以及使用新Objective-C消息和存根合成语法合成存根。

总结

在本节课中，我们一起学习了如何利用LLVM ORC JIT来动态加载和运行原本为静态链接设计的程序，从而为Xcode的Swift UI预览功能提供动力。我们探讨了覆盖插件如何实现快速函数替换，内存管理器如何支持代码签名以兼容macOS强化运行时，以及通过动态加载器集成使JIT代码对预编译代码可见的核心创新。我们还了解了为处理大规模、复杂程序所做的各项性能优化，以及支持各种特殊构建配置和链接选项的改进。

ORC JIT现在能够加载原本用于静态链接的程序，可以扩展到非平凡的程序规模，并支持不寻常的构建配置和执行环境。借助动态加载器支持，预编译代码可以与JIT编译的代码交互，就像JIT编译的代码是预编译的一样。我们期望这个项目的许多好处能够惠及LLVM JIT的其他客户端，并认为这里存在新的开发者工作流机会有待探索。

021：缓解侧信道漏洞的挑战

概述

在本节课中，我们将要学习在LLVM编译器中缓解侧信道漏洞所面临的挑战。我们将探讨侧信道攻击的原理、为何传统的“恒定时间”编程模型不足，以及如何通过编译器层面的代码转换来应对这些威胁。

背景：安全关键代码与侧信道

安全人员主要关注处理秘密数据的软件，我们称之为安全关键代码。例如，密码哈希比较和加密质量检查都属于此类代码。我们并不关心通用代码，只专注于强化这些安全关键的加密代码。

侧信道攻击通过程序产生的非预期副作用来泄露数据。例如，程序可能将数据写入缓存，攻击者通过观察缓存状态来间接读取数据。另一种常见的类型是时序侧信道，例如，一个程序在密码正确时运行得更快，攻击者可以利用这种时间差异来提取密码。

恒定时间编程及其局限性

为了应对这些威胁，恒定时间编程一直存在。顾名思义，恒定时间编程意味着编写的程序对于任何输入，其执行都不会产生可变的副作用。例如，代码中应避免基于秘密数据的分支，或者在遍历列表时，即使找到了目标也不提前退出。

然而，恒定时间编程足够吗？不幸的是，答案是否定的。我们已经看到，像Spectre和Meltdown这样的攻击也能破坏恒定时间程序。

不仅如此，即使一个程序现在是完美的恒定时间，仍然存在许多微架构侧信道，它们源自缓存、预取器等，仍能产生攻击者可利用的显著差异。

微架构侧信道

假设有一段恒定时间代码，它看起来应该以恒定时间执行。但硬件内部可能会决定进行一些优化，这些微架构优化可能导致程序的执行产生不同的副作用。最常见的例子之一是推测执行。

推测执行很简单，我们看到一个条件判断，然后是一个移动操作。实际上，程序中的条件是“如果索引在边界内，则执行此移动”。但CPU决定推测性地执行移动操作，当它发现条件不成立时再回滚。然而，在这个推测性或瞬态执行状态期间，缓存状态已经被改变，这可能帮助攻击者提取信息。我们已经看到Spectre攻击有多么严重。

但这是一个复杂的例子。让我们看一个非常简单的例子，也是我最喜欢的，叫做静默存储。假设有两个存储操作指向完全相同的地址，但它们来自不同的寄存器。在硬件中，如果这两个值（例如EAX和EBX的值）相同，那么第二个存储操作就会被静默化，直接被优化掉。这在硬件层面很有道理，我们不想浪费宝贵的周期再次执行存储操作。然而，从安全角度看，这很糟糕。攻击者已经证明，这种看似微不足道的优化可以被利用。

类似地，一个更简单的例子是加法指令。当其中一个寄存器或参数为0时，我们期望它不应该进入ALU或任何复杂的电路，而应该直接返回第二个参数，因为 a + 0 = a。但这可能产生时间差异。因此，当你的一个变量为0时，攻击者会知道这个变量是0，因为加法指令返回得更快。

这些微架构优化是必要的，我们无法摆脱它们，但它们可能产生攻击者可观察到的副作用，从而导致数据泄露。最大的启示是，这些优化会破坏你的恒定时间代码。所有投入编写恒定时间加密库和程序的努力都可能被编译器或硬件优化所抵消。

编译器缓解策略：选择性代码转换

我们意识到，程序中存在一些代码模式。作为编译器工程师，我们非常关注这些模式，并将它们转换为更高效、更优化的版本。但同时，这些代码模式也可能被硬件用来触发那些可能易受攻击的优化。

那么，我们能否直接关闭那些不好的优化呢？我们不能关闭推测执行，那会导致巨大的性能损失。我们能否为关心安全的人提供一个开关呢？目前还没有人发布这样的功能。我们能否以编程方式关闭它呢？例如，在执行加密代码时关闭优化，在非关键部分再正常开启？这在一定程度上是可能的，但支持非常有限。

因此，我们思考的另一面是：能否消除所有不应触发那些优化的代码模式？例如，如果我们知道程序的某个特定区域处理秘密数据，能否转换该部分程序，使其在硬件上执行时永远不会触发那些优化？这很聪明，但也很复杂，因为我们必须覆盖所有可能的情况。然而，事实证明这是一个很好的备用和兜底解决方案，因为你可以在“第0天”就实施，也可以为旧硬件提供保护。

这种“第0天”解决方案意味着，一旦出现新攻击，你就可以在编译器中推出这种缓解措施。所有用该编译器编译的代码都不会调用那个可能易受攻击的优化。这种方法也适用于你的旧硬件。

我们需要对程序进行污点分析，以找出程序中处理秘密数据的区域，这些是我们需要保护、不应触发那些优化的区域，然后对它们进行转换。

最简单的案例是Spectre缓解，大多数人可能都熟悉，即插入屏障指令来阻止推测执行。这是相同的概念，只是我们已经在实践中使用了它。假设有一段易受攻击的代码，并且这段代码位于安全关键区域，我们插入一个屏障。由于这个屏障，我们知道这个移动操作不会被推测执行。这实际上是在某种程度上关闭了优化。如果我们能在正确的时间、正确的地点、针对正确的数据（即我们的秘密）这样做，我们就能有效地保护秘密免受Spectre攻击，而无需承受完全关闭推测执行或为程序中所有操作插入屏障所带来的性能损失。

类似地，对于静默存储，你有两个移动操作。如果我在它们之间再插入一个移动操作呢？我插入另一个存储操作，确保R11中的值永远不会等于EAX和EBX。我们实际做的是，取EAX和EBX的值，将它们拆分。我们从EAX取高半部分，从EBX取低半部分，然后将它们反转。这样做，我们保证生成的值永远不会等于EAX或EBX。这使得静默存储不会因为EBX而触发，因为R11永远不会等于EBX。同时，它也不会因为R11本身而触发，因为R11也永远不会等于EAX。

我希望以上内容能让大家对这些攻击的工作原理，以及我们如何通过选择性代码转换在编译器中禁用或缓解它们，有一些直观的理解。

实施缓解措施面临的挑战

现在，让我谈谈我们在实施这些缓解措施时面临的一些挑战。过去两三年，我一直在为不同复杂度的攻击实施这些措施。

第一个挑战是，我们应该在哪里实施这些转换？我们是在LLVM IR层面、后端寄存器分配之前，还是之后？假设你关心所有存储操作（如静默存储），我们可能希望在寄存器分配之后进行转换，因为那时溢出和填充已经发生，我们知道所有可能发出的新存储操作。但对于某些缓解措施（例如分配影子内存），那个抽象层次可能太低了。你可能希望在LLVM IR层面做，而不是在寄存器分配之后的MI层面，那会非常痛苦。我们意识到没有完美的解决方案。对我们来说，最好的方法是：在LLVM IR层面能做的就在LLVM IR做，然后向下传递信息，例如标记这些数据的来源。

另一个有趣的问题是，如果我们想在寄存器分配之后实施静默存储转换，我们需要预留R11寄存器。这是一个非常常见的做法。在Helium代码中，我们必须为x86手动预留R11，目前没有更好的方法。但这些改变甚至会影响编译周期的行为方式。

下一个挑战是，我应该在何处调度我的缓解通道？我想把它放在最后，因为我不希望任何其他优化或缓解措施撤销它。但如果每个人都想放在最后，那“最后”是什么？如果你有两个缓解通道，它们之间的顺序如何？考虑到新威胁不断涌现，以及内存安全可能在下一个十年内得到解决，侧信道可能会成为更大的挑战，我们需要处理越来越多的缓解通道。

还有一个视角的转变。当我们进行优化时，我们只关心两件事：等价性和速度。我们找出那些我们有100%把握能保持等价且更快的“低垂果实”。但对于缓解措施，你不能遗漏任何东西。如果你遗漏了一个代码模式，那就是一种虚假的安全感，一个潜在的漏洞。而且这是一个更难发现的漏洞，因为你声称这是一个已缓解的二进制文件。

一些开放的研究问题包括：作为编写这个通道的人，我如何推理排在我前面的通道？这也与编写恒定时间程序的人所面临的痛苦相关，他们希望使用编译器，希望他们的恒定时间程序运行得快，但编译器是如此复杂的基础设施，其行为方式直到完成前都不容易理解。

另一个挑战是内联汇编。编译器中的内联汇编只是一个字符串，我们如何处理它？如果你有一个已缓解的二进制文件，但没有处理内联汇编，那效果就不好，因为攻击者仍然可以找到漏洞。在我们的实验中，几乎所有加密库都有一个没有内联汇编的版本（即非高度优化的版本），我们使用那个版本。但最好的方法是在二进制层面（如二进制重写）处理，但这也有其自身的挑战，因为你无法带入所有分析结果。

此外，还有一个视角：为了弄清楚两个转换是否冲突，或者如何排序它们，是否需要为每个转换定义语义？这需要大量的工作。但我们能否有简单的后转换检查器呢？例如，我生成了这个二进制文件，它是否还包含这个代码模式？这些是可以用二进制写入器实现的更简单的事情，我们实际上也在这么做。所以，一旦我实现了缓解措施，我会获取二进制文件，编写一个检查器，然后检查是否有遗漏。令人惊讶的是，很多时候我以为自己做得很好，结果却发现遗漏了很多代码模式和情况。编译器是如此复杂，后端也是如此复杂。即使你想在寄存器分配后写一个MI通道来支持所有存储操作，但要支持所有生成存储的地方（例如，调用约定中的大小等），也是非常痛苦的。

总结与展望

在本节课中，我们一起学习了侧信道漏洞的挑战以及编译器层面的缓解策略。

我想在演讲结束时让大家记住以下几点：基于编译器的防御措施对于应对这些新兴威胁非常灵活且及时。明天出现新攻击，你可以更快地推出基于编译器的缓解措施。我们在Spectre事件中也看到了这一点，基于编译器的防御措施出现得非常快。

我100%同意这些措施可能更慢的观点，但我也相信它们是必要的，因为你无法回头更改硬件。例如，我使用一台M1芯片的笔记本电脑。如果我关心M1上的某个攻击，我可能必须换到其他修复了该问题的新硬件，或者，也许这些编译器转换可以帮到我，因为我只需要将它们应用于实际处理密码或私钥的加密代码。从更大的视角来看，这部分代码相对于我们使用的总代码量来说可能很小，其开销可能很大。但从整体视角看，这部分小代码上的巨大开销是可以管理的。

编译器是惊人的、优秀的基础设施，非常适合进行优化、转换和分析。但它们最初并不是为缓解措施而设计的，这就是为什么我们面临这些挑战。这并不是说编译器不好，它们只是不为缓解措施而生。但仍有改进编译基础设施的空间，使得缓解措施也能成为其中的一部分。

问答环节

提问者：感谢精彩的演讲，发人深省。我有一个问题：你如何评估你的程序转换是否有效？特别是，例如对于静默存储优化，如果在一块硬件上足够欺骗微架构优化，但在另一块可能已经存在的新硬件上无效，怎么办？你如何验证这类事情？另外，随着新硬件的出现，你如何检查某些缓解措施不会使其他缓解措施退化？这里的评估策略是什么？

回答：正确。我们实现的东西是非常特定于架构和微架构的。例如，我们为x86的某个特定微架构实现的东西，就只适用于那个。我们并不声称它适用于任何其他版本或新出现的版本。有趣的是，我们的很多工作也依赖于攻击者社区，他们找出某个特定微架构对这些攻击是脆弱的，并且他们能很好地给出攻击信息。可以说，目前这些转换非常简单，如果你看到新攻击的规范，你可能会自己得到答案。但目前，我们并不那样做。

本节课中，我们一起学习了：

  1. 侧信道攻击的基本原理：攻击通过程序的非预期副作用（如缓存状态、执行时间）泄露秘密数据。
  2. 恒定时间编程的不足：仅靠软件层面的恒定时间编程无法抵御由硬件微架构优化（如推测执行、静默存储）引入的侧信道。
  3. 编译器缓解策略的核心思想：通过选择性代码转换，消除程序中可能触发有害硬件优化的代码模式，从而在秘密数据处理区域“关闭”这些优化。
  4. 实施中的主要挑战：包括转换实施的位置（IR vs. 后端）、通道调度与顺序、确保覆盖无遗漏、处理内联汇编以及验证转换有效性等。
  5. 编译器防御的优势与必要性：提供了灵活、及时的响应手段，尤其对于无法立即更新硬件的旧系统至关重要，尽管可能带来性能开销，但在安全关键代码上是可管理的。

编译器是强大的基础设施，虽然最初并非为安全缓解而设计，但通过持续改进，它将成为应对日益复杂的侧信道威胁的关键组成部分。

022：自动化编译器工程中的搜索环节

在本节课中，我们将学习如何将编译器工程中的部分工作自动化。具体来说，我们将探讨如何通过扩展MLIR的Transform方言，将优化流水线中的参数调整和方案选择等“搜索”环节表达为IR，并利用约束求解器自动完成组合搜索，从而将工程师从繁琐的手动调优中解放出来。

概述：编译器工程中的日常挑战

在我看来，编译器工程涉及大量流水线和参数（旋钮）的调整工作。一个典型的工作日可能是这样的：我们有一个现有的流水线需要优化。凭借经验和洞察力，我们决定在特定位置引入新的优化以提升特定工作负载的性能。然后，我们从众多可能的优化中，首先尝试优化F，并评估其性能。如果性能不达标，我们就开始调整F的所有可用参数，并反复评估。如果F不行，我们再尝试优化G，并重复调整和评估的过程。有时我们甚至会尝试不同的优化顺序，例如先F后G，或者先G后F。这个过程包含了大量手动工作。

这里存在一个潜在的改进方向：如果我们能够将这类优化计划表达出来，或许就能将这个过程机械化、自动化。

课程结构

在本课程中，我们将按以下步骤展开：

  1. 快速回顾调度（Schedules）及其与Transform方言的关系。
  2. 介绍用于调优参数的IR，并给出这些操作的非确定性语义。
  3. 引入新的操作来处理可选的变换序列。
  4. 探讨如何优化这个“优化循环”本身。

1：什么是调度？

程序实际上由两部分组成：算法（Algorithm）和调度（Schedule）。算法定义了功能行为，即输入如何映射到输出。而调度则规定了执行应如何进行，例如使用哪些内存、采用何种循环顺序等。

让我们通过Halide的一个著名例子来具体说明。我们有一个对图像应用模糊效果的高级算法。而在特定硬件上运行的实际程序，看起来更像右边这样。在传统编程范式中，程序员最终需要编写这个结合了算法功能行为和所有性能优化细节的低级程序。

根据调度优先的范式，我们可以采用不同的方式：我们可以编写算法的高级代码（功能行为），并声明一个调度，来指导如何对算法应用何种优化和降低，以得到优化后的低级代码。因此，程序员不直接编写低级程序，而是编写算法和一个描述如何优化的调度，然后由某种机制将调度应用到算法上。

调度优先范式有许多实例，包括TVM和Halide，它们都共享相同的基本理念：我们希望对代码变换进行声明式描述。

2：MLIR的Transform方言

这引出了MLIR的Transform方言，它只是调度范式在MLIR中的另一个实例。MLIR的Transform方言允许你编写IR来描述应对其他IR进行何种变换。因此，调度就是Transform IR，它使用Transform方言来指导如何修改有效载荷IR（Payload IR），即我们想要降低和优化的常规IR。

例如，我们有一些linalg.generic操作，调度指示我们找到任何具有一个并行迭代器的操作（即找到这个linalg.generic），然后对其应用循环分块（tiling）。我们将有效载荷IR和调度提供给一个称为Transform解释器的机制，该解释器应用调度，最终得到我们想要的、带有额外循环的IR。

更多背景信息：对于不熟悉MLIR的人，我们可以逐步降低（lower）各种方言。我们从非常高级的表示开始，然后转到linalg，接着可能添加一些分块，最后可能通过分解来消除linalg。Transform方言的理念是，这些步骤中的每一步都可以描述为一个小调度，我们甚至可以复用已有的Pass。因为Transform方言就像一种合适的编程语言，这些调度也可以组合，因此我们可以用一个大的调度来描述整个降低过程。这本质上替代了你的优化流水线。

3：引入参数选择到调度中

一旦你开始尝试这些调度或流水线，很快就会发现到处都是“魔法值”。例如，调度告诉我们要找到MatMul并对其应用分块，我们可能需要进行缓存分块和寄存器分块，但我们并不确切知道应该使用哪些值。对于每个想要使用的不同分块尺寸，我们最终都会得到一个不同的调度。这通常涉及手动调整或通过脚本生成。无论哪种方式，都很难追踪我们愿意考虑的不同选项。我们需要一个更上层的东西来讨论这些不同的选项。

我的提议是，我们可以将这种“选择”的概念具体化到调度本身中。我们这样做：对于确定M维度缓存分块的大小，我们有一个选择。我们通过引入一个transform.tune操作（附带一个描述性名称）来表达这个选择。我们说它有许多不同的选项。这个操作的语义是，它产生一个SSA值，该值的实际值是选项之一（例如16、32或64）。我们可以有多个这样的独立节点。

我们首先观察到，我们可以直接遍历IR，提取出这些引入的参数所构成的参数空间。但这带来了一个问题：这些transform.tune旋钮根据Transform解释器是不可执行的，因为这里存在非确定性（值可能是16、32或64之一），解释器不知道如何处理这种非确定性。所以我们需要帮助它，我们需要做点什么。

4：映射到约束求解问题

我们能做的第一件事是提取那个约束问题（搜索空间）。然后，我们需要一种机制来在这些不同的值/选项之间做出选择。之后，我们可以有某种驱动程序或脚本，自动将带有选择的调度重写为没有选择的调度。这里，我们只是将调优旋钮映射为常量。这样，我们就从一个不可执行的调度变成了一个可以执行的调度。

我们通过SMT（可满足性模理论）方言来处理这些约束问题。SMT本质上是一种用于表达约束问题的语言或机制。细节并不复杂，但我会给你一个概览：在左边，我们有一个transform.tune旋钮，我们将把它映射到SMT。这意味着我们将旋钮解释为一个可以取不同值的变量。这里我们声明有一个变量（恰好是这个名称），并且我们说它可以是16、32或64。因为我们现在仍在MLIR IR中，我们可以将其转换为SMT求解器使用的语言，这是一个非常直接的映射。此时，我们可以将其交给任何现成的SMT求解器软件。

SMT求解器软件可以枚举MC变量的所有有效赋值。然后，我们可以使用这些赋值将非确定性调度重写为确定性调度。在这个简单例子中，我们最终从一个非确定性调度得到三个确定性调度。这是一个相对直接的方法，适用于我们只有这些独立旋钮的情况。枚举独立旋钮的所有赋值并不复杂。

5：表达旋钮间的约束

我们希望能够在表达旋钮的同时，直接表达我们对于不同可调参数有效赋值的约束。因为我们知道我们将映射到SMT方言，我们可以利用MLIR中方言混合的特性。例如，对于寄存器分块，我们之前说只有这些有效组合。但实际上，我们知道约束更像是这样：reg_tile_m * reg_tile_n <= 64。

现在，我们可以用多个旋钮来表达它，并表达实际的约束。我们引入另一个新操作，它表示：我知道这些transform参数实际上将被视为SMT变量，所以让我们将每个参数映射到一个SMT变量，然后我就可以直接使用SMT方言来表达我想要的任何约束。其转换是直接的：我们独立转换每个旋钮，然后约束只是我们在特定区域中内容的副本。我们将其映射到SMT，要求SMT软件枚举所有有效赋值，我们得到许多赋值，所有这些都满足我们设定的特定约束。

6：处理变换序列的替代选择

我认为我们需要更进一步，以解决我们最初提出的例子。在优化调度时，你会发现你会说：我要么做这个变换，要么做那个变换。我希望能够表达这种选择。显然，仅用参数调优是无法做到这一点的。参数调优适用于调整变换或Pass的参数。而这里，我们要表达的是在Pass或变换之间的选择。

这实际上很容易编码。我们引入一个新操作transform.tune.alternatives，它表示我们将使用第一个区域或第二个区域。这个操作的语义是：用其区域之一替换这个操作。因此，一旦有了这个操作，我们本质上就得到了我们之前想要的两个独立调度的含义。

为了将其映射到SMT，transform.tune.alternatives可以有任意数量的区域。我们需要表示选择的是：对于这个特定的tune.alternatives，我们有一个SMT变量，其取值范围是从0到n（n+1个选项）。然后我们断言情况确实如此。接着，我们使用公式来检查子变量实际上被赋值为0、1或n。这允许我们检查选择实际上走了哪条路。然后，我们需要确保在这里面的任何SMT操作、任何旋钮或其他选择也被转换，并且这必须取决于那个特定选择是否成立。

7：表达选择间的依赖关系

对我来说，这仍然不够。通常我们不仅关心在调度中特定位置做出选择，后面的选择可能依赖于前面做出的选择。例如，在左边，我先进行并行分块，然后稍后我想将scf.forall操作替换为scf.parallel操作，这显然依赖于之前进行了tile_with_forall。在右边，我先进行顺序分块，然后我会对顺序循环做一些操作。

最直接的做法是将此表达为两个独立的选择：首先是tile_with_forall和tile_with_for之间的选择，然后是映射到并行或进行coalesce之间的另一个独立选择。但这样我们最终会得到四个调度，因为第一个有两个选择，第二个有两个选择。

为了使第二个选择依赖于第一个，我们引入一个标记（marker），它是第一个tune.alternatives的结果，允许我们跟踪选择的方向。在这个例子中，如果我们使用并行区域，我们将标记设置为marked_parallel（值为1）；如果我们使用顺序区域，则标记为marked_sequential（值为0）。然后，对于第二个依赖于第一个的替代选择，我们可以直接使用Transform方言自身的操作来断言标记应该等于marked_parallel。当我们对并行循环进行操作时，最好确保标记表明我们之前对并行循环进行了操作。对于要对顺序循环进行操作的情况，我们也断言最好确保我们之前对顺序循环进行了操作。

然后我们将其降低到SMT，这相当详细。我们将其提供给SMT求解器软件，在这种情况下，只有两个满足所有约束的变量有效赋值。因此，我们能够用两个替代选择编码出我们只愿意考虑的两个具体有效调度。

8：编码完整的优化计划

现在，我们拥有了将所有工具用于将我们“平均工作日”的平均计划编码到IR中。在左边，我们有想要优化的流水线，我们已经在IR中看到了它的轮廓。现在我们要说：我们要么做F，要么不做F。对于F，我们有一些旋钮，我们可以直接在调度中添加这些旋钮。如果我们不做F，那么我们就做G。G也有其旋钮。现在我要表达另一个选项：我先做G，后做F。当然，它也有其操作。所以，我要么在这里做G，要么什么都不做。使用一个标记来记录发生了什么。然后，我将利用这些tune.alternatives选择来表达“F之后的G”，仅当G尚未执行时才执行。否则，我们什么都不做。如果我们没有这一行，就可能出现先G、后F、再G的情况。现在，我们对愿意考虑的具体调度版本表达了大量的控制。

总结与自动化流程

总结一下，现在我们拥有了自动化调度/流水线优化的完整机制：

  1. 我们拥有带有旋钮和替代选择的调度。
  2. 我们将其映射到SMT输入（SMT-lib），这是SMT求解器可以处理的标准格式。
  3. 我们将其交给约束求解器（SMT求解器），确保它只挑选满足所有约束的有效赋值。
  4. 然后，我们可以将调度重写为没有旋钮和替代选择的具体调度。
  5. 我们可以将这些具体调度与想要优化的有效载荷一起提供给Transform解释器。
  6. 我们拥有用于评估和评分的测试工具，我们可以循环运行此过程，例如自动收集评分最高的旋钮和替代选择的赋值。

目前我们已经实现了这一部分。我们还想更进一步，例如优化约束求解过程，使其搜索更智能；或者同时优化许多不同的有效载荷，因为我们可能希望针对不同的工作负载优化整个流水线。

课程总结

本节课我们一起学习了：

  • 我们引入了一种将参数选择具体化到IR中的方法，这为我们提供了具有非确定性语义的旋钮操作。
  • 我们使得讨论变换序列的替代选择成为可能，这超越了常规的自动调优。
  • 关键点在于，通过联合约束，我们可以对我们愿意考虑的不同具体调度表达大量的控制。如果我们不能做到这一点，很容易陷入难解性问题。
  • 我们可以将搜索满足约束的赋值的任务交给求解器，从而自动提取可执行的调度。

023：过去与未来

在本节课中，我们将回顾 LLVM Libc 项目自2019年提案以来的发展历程，探讨其当前用户需求，并介绍一个面向未来的全新设计提案。我们将重点关注项目的模块化、跨平台支持和社区导向，并规划实现一个具体目标：在2026年底前，让 LLVM Libc 能够支持 Clang 编译器，使其达到生产就绪状态。

项目起源与2019年提案

首先，让我们回顾一下 LLVM Libc 项目的起点。2019年，Shiva Chandra 撰写了一份提案，为 LLVM Libc 确立了五项指导原则。

以下是该提案的核心原则：

  1. 作为库构建：Libc 应真正遵循“库”的哲学，而非一个难以分解的单体。
  2. 支持静态链接：优先支持将库静态链接到应用程序中。
  3. 将标准视为指南：以满足用户实际需求为目标，而非刻板遵循标准。
  4. 谨慎对待供应商扩展：对非标准功能持审慎态度。
  5. 成为 LLVM 工具链的典范：充分利用 LLVM 生态的各种工具来构建高质量的库。

关于“作为库构建”这一点，原提案中有这样一段描述：“尽管 C 标准库名义上是一个库，但大多数实现在实践中是相当单一的整体。” 这就像一块紧密拼接的乐高积木，难以单独替换其中的某一块。

该提案也明确了一些非重点领域，例如动态加载和增加更多架构支持。不过在实际发展中，我们后来增加了对多个新架构的支持。

当前进展与用户生态

上一节我们回顾了项目的起点，现在来看看过去六年我们取得了哪些成果。

  • 代码规模：我们新增了超过 1000 个函数，其中 480 个是数学函数。
  • 架构支持：目前支持大约 6 种处理器架构（具体计数方式不同，结果在4到8之间）。
  • 社区贡献：约有 50 位贡献者提交了超过 10 次提交，总计超过 220 人至少有一次提交。我们目前拥有 11 位维护者。

那么，是谁在使用我们构建的这个库呢？主要有以下几类用户：

  • 部署在已知环境：
    • 容器化服务器（如 Google 所用）。
    • 嵌入式设备（如新款 Pixel Buds）。
    • 甚至可以为 GPU 构建（Joseph Huber 的工作）。
  • 作为源代码库使用：
    • LLVM 子项目通过 libc-utils 项目使用 Libc 源码，例如 libc++ 用于字符串到浮点数的转换，OpenMP 用于 printf 内部功能。
  • 外部用户集成：
    • Android 的 Bionic Libc 引入了我们的宽字符函数。
    • Fuchsia 正在用新的 LLVM Libc 实现替换其部分 Libc 函数。
    • musl 提供了一个实验性配置，允许使用 LLVM Libc 替代原有的 C 运行时。

用户需求与设计优先级

了解了用户群体后，我们来看看他们具体需要什么，以及这对我们的设计意味着什么。

用户通常采用“在应用中构建 Libc”的模式。这意味着他们使用静态链接，或者像 Google 的“运行时按需链接”那样，将库和应用程序一起构建和链接。这种模式的优势在于，它避免了传统动态链接库模型中，应用程序需要同时兼容新旧版本库 ABI 的难题。只要应用程序的二进制接口稳定，就可以直接部署包含新版本库的完整新应用。

除了构建模式，用户还有以下核心需求：

  • 可移植性：包括易于移植到新架构（向上可移植），以及在不同硬件上提供一致的 API（向下可移植）。
  • 质量属性：性能、代码大小和数学精度都是接口的一部分。不同的应用场景需要在三者之间做出不同的权衡，因此库需要一定程度的可定制性。

相应地，以下是我们当前的非重点领域：

  • 动态加载：由于用户倾向于静态链接，动态加载不是当前焦点。
  • ABI 稳定性：对于静态链接或源码集成的用户，只要他们的应用能重新构建，库的 ABI 是否稳定并不关键。
  • 遗留的有问题特性：例如将区域设置（locale）作为隐式参数（如 printf 依赖全局区域设置），这种设计在现代环境下容易引发问题，更好的方式是使用显式参数。

新的设计原则：模块化、跨平台、社区导向

基于对用户需求的分析，我提出以下三项新的指导原则，它们将指引 LLVM Libc 的未来发展。

1. 模块化
我们希望函数之间相互独立，实现“垂直模块化”，以便能够单独部署某个函数，而不必引入一堆不需要的依赖。同时，我们希望操作系统层的接口是通用的，实现“水平模块化”，这样在移植到不同平台（如 Linux 到 Windows）时，可以轻松替换底层的 OS 模块，而无需重写上层逻辑。

2. 跨平台
为了实现跨平台，我们应尽可能使用 C++ 编写可移植代码，并自动生成汇编，以减少需要手动维护的平台特定汇编代码。我们需要清晰地划分平台无关代码、平台通用代码和平台特定代码的层次。目标是实现“一个前端，多个后端”——相同的 API，通过不同的后端适配到各个系统。

3. 社区导向
一个优秀的库离不开一个健康、成长的社区。我们需要营造一个欢迎新人、友好、尊重并能建设性处理分歧的环境。社区应为成员提供新的机会，让他们能够成长并承担更重要的角色。

具体目标与实施路线图

前面我们讨论了新的设计原则，现在让我们聚焦于一个具体、可衡量的目标。

我们计划在 2026 年底前，让 LLVM Libc 达到能够支持 Clang 编译器的生产就绪状态。选择 Clang 作为目标是因为它是一个真实、行为良好、广泛可用的生产级程序。实现这一目标也将意味着 LLVM 工具链实现了自举（self-hosting），这是一个令人兴奋的里程碑。

以下是实现该目标的路线图：

2025年底前：

  • 完成本战略的宣讲。
  • 完成所有必要的代码清理和重构的设计工作。

2026年初：

  • 实施上述设计好的清理和重构工作。
  • 实现 Clang 所需的一系列函数，包括宽字符文件函数、部分 pthreads 功能以及一些 POSIX 系统调用包装器。

2026年底前：

  • 建立 Clang/LLVM Libc 的持续集成构建机器人，确保 Clang 不仅能构建，还能通过所有测试。
  • 进行最终打磨，使 Clang + LLVM Libc 达到真正的生产就绪水平。
  • 不仅实现 Clang 所需的函数，而是完成相关函数集的全部实现，避免库中出现奇怪的缺口。
  • 设定下一个激动人心的项目目标。

问答环节摘要

在演讲后的问答环节，讨论了一些关键问题：

  • 如何替换特定函数（如 malloc）？ 主要通过构建系统实现。项目目录结构支持按平台/架构提供特定实现，构建时会选择最具体的版本。
  • 静态链接的代码大小影响？ 预计影响不显著。对于嵌入式设备，死代码消除会移除未用部分；对于服务器，静态链接增加的大小上限大约相当于动态库文件的大小（如 glibc 的 .so 文件约 2MB）。
  • 静态链接下的安全更新？ 在容器化部署等场景中，通常需要重新部署整个应用程序，因此可以在新版本中直接包含修复后的库。
  • 对动态加载和插件支持的计划？ 最终希望将动态加载作为一个可选项。初步阶段，可以尝试配置 Clang 禁用插件加载，或按需构建动态加载支持。
  • 对 Windows/macOS 的平台支持？ 目前不完整，主要缺乏相关贡献者。重构以改进 OS 抽象层后，移植工作会更容易。
  • 将 LLVM Libc 作为库集成到其他项目的难易度？ 目前已经比较容易，许多代码（如 printf、字符串转浮点数、数学函数）已经是平台无关的，可以方便地抽取使用。
  • 如何开始贡献？ 可以访问 libc.llvm.org，参加每四周一次的公开会议，或通过 Discourse、Discord 频道交流。

总结

本节课中，我们一起回顾了 LLVM Libc 从2019年提案至今的发展，分析了其用户群体和核心需求。我们重点介绍了一个面向未来的新设计提案，其核心是模块化、跨平台和社区导向三项原则。最后，我们设定了一个清晰的目标：在2026年底前，让 LLVM Libc 能够支持 Clang 并达到生产就绪状态。这是一个充满挑战但令人兴奋的旅程，欢迎社区的每一位成员参与其中，共同攀登 LLVM Libc 的完整阶梯。

024：优化Flang的优化器

在本教程中，我们将学习如何分析和优化Flang编译器（LLVM的Fortran前端）的性能。我们将通过分析SPEC CPU 2017基准测试中的具体案例，探讨如何识别性能瓶颈并实施优化，最终目标是使Flang的性能与GCC的差距控制在10%以内。

1️⃣：Flang简介与性能评估

Flang是LLVM项目的Fortran语言前端。它的工作流程是将Fortran源代码逐步降低（lower）为不同的中间表示（IR）。首先，Fortran代码被转换为高级Fortran IR（HLFIR），然后进一步转换为FIR（Flang IR），最终转换为其他MLIR方言和LLVM IR。从这一点开始，后续的编译步骤由MLIR和LLVM完成，最终生成可执行文件。

在从Fortran到HLFIR的降低过程中，除了HLFIR操作外，还会使用其他MLIR方言，如func和arith。HLFIR保留了Fortran语言结构和语义的许多细节，因此它是执行Fortran特定优化的理想阶段。

在尝试优化Flang之前，首先需要评估其性能。我们使用SPEC CPU 2017基准测试中仅包含Fortran代码的部分。测试运行在一台Nerverse VI机器上。我们的目标是，在任何基准测试上，Flang的性能不应比GCC 14低10%以上。

截至今年4月，LLVM和GCC的最新稳定版本分别是20和14。性能对比图表显示，在大多数基准测试中，Flang的性能与GCC 14相当或更好。然而，在cactuBSSN（简称c4）测试中，Flang的性能低了18%。因此，c4成为了一个很好的优化候选对象。

2️⃣：识别性能瓶颈

为了识别占用最多执行时间的函数，我们在Ubuntu机器上使用了perf工具。分析发现，性能问题分散在整个程序中，这使得识别差异最大的函数变得有些困难。

以下是perf report输出的c4测试中最热的10个函数列表。我们选择了array_props_s_w函数，因为其样本数几乎是GCC版本的两倍，表明性能差距显著。

然而，在尝试检查该函数时，只列出了汇编指令，没有对应的Fortran源代码。这是因为与GCC不同，当启用LTO（链接时优化）时，LLVM有时会丢失调试信息。我们的解决方法是禁用LTO，以便能够将汇编指令映射回Fortran结构。

禁用LTO后，我们现在可以看到与汇编指令并列的Fortran源代码。array_props_s_w是一个大型函数，但引起我们注意的是其中三维数组的初始化部分。可以看到，LLVM使用了多个嵌套循环结构，而GCC仅使用两个简单的循环。问题似乎在于LLVM没有意识到这是对整个数组的连续初始化，因此未能使用单个循环来优化。

3️⃣：数组初始化优化

为了探究问题的根源，我们检查了代码生成的每个阶段，从HLFIR开始。在HLFIR中，整个初始化由一个单独的hlfir.assign操作完成。当它被转换为FIR时，变成了三个嵌套循环，每个数组元素都通过store double指令初始化。在LLVM IR中，这又变成了一系列getelementptr指令后跟store double指令。

这种情况的发生是因为两个内层循环被完全展开了。但由于数组在内存中是连续的，也许我们可以让Flang将assign操作降低为使用单个循环，这可能有助于LLVM生成更优化的LLVM IR。

要实现这一点，首先需要对多维数组进行线性化，以便能够用单个循环遍历它。以下是执行此优化的核心代码逻辑，但我们将重点关注应用此优化后生成的FIR代码。

可以看到，三个嵌套循环被替换为一个单循环，该循环遍历扁平化的数组，为每个元素存储double值。在LLVM IR中，现在也变成了一个单循环，并且这次甚至能够将store double指令成对分组。最终，在汇编指令中也可以确认，初始化部分只剩下一个循环。

然而，不幸的是，最终c4的执行时间几乎没有变化。可能的原因包括：array_props_s_w是一个大型函数，数组初始化只占其运行时间的一小部分；此外，该函数中还存在其他分散的性能问题。不过，在其他测试程序中，我们观察到了显著的性能提升。如下表所示，对于64kB到256kB大小的数组，速度提升了约30%。另一个好处是，数组初始化所需的指令数在c4案例中从38条减少到了8条。

4️⃣：字符串比较优化

这是我第一次尝试优化c4。一段时间后，我再次尝试优化它。那时GCC 15已经发布，我使用了LLVM的主干版本。同时，由于其他Flang开发者的贡献，array_props_s_w函数的性能差距已经缩小。

我使用火焰图来帮助识别在比较GCC和LLVM时差异最大的函数。识别出的函数是twamax_is。对比火焰图可以看到，LLVM编译版本在该函数上花费的时间显著高于GCC版本。放大观察，我们发现LLVM版本调用了trim和free函数，而GCC版本只调用了一个不同的字符串比较运行时函数。

perf输出证实了这一点：LLVM进行了两次trim调用、两次free调用，并且字符串比较是内联进行的。这个函数实际上很简单，只是比较两个修剪（trim）后的字符串。有趣的是，GCC能够简单地消除这两个trim调用。

让我们看看这是如何实现的。首先需要记住，在Fortran中，字符字符串在右侧用空格填充。此外，查阅Fortran语言规范可知，在比较操作中，较短的字符串操作数总是用空格填充以匹配另一个操作数的长度。所有这些意味着，可以安全地去掉trim调用，因为关系运算符本身已经考虑了右侧的空格。

为了实现这种表达式简化，第一步是添加一个新的HLFIR操作来表示对trim的调用。这使得更容易识别并在可能时移除该调用；当无法移除时，该操作会被转换为一个行为与之前相同的运行时调用。

在此基础上，我编写了一个Pass来简化字符比较。其核心可以总结为以下表达式：两个修剪后字符串的比较，变成了这两个字符串本身的比较。

以下是执行此转换的核心代码，但我们将跳过代码细节，重点关注优化应用前后的HLFIR代码。可以看到，在优化后的版本中，两个chartrim操作以及相应的destroy操作都被移除了。

这项优化的结果是，c4的执行时间从217秒下降到280秒，估计带来了约3.5%的速度提升。

5️⃣：其他贡献与未来计划

我想提一下，其他开发者也为Flang的性能提升做出了许多贡献。本幻灯片列出了一些近期的贡献。在当前LLVM主干版本与GCC 15的性能对比图表中，红色柱代表LLVM（从20版到主干版），蓝色柱代表GCC 14和15。可以看到，现在唯一一个Flang性能比GCC 15低10%以上的基准测试是exchange2，差距为11%。尽管c4仍有改进空间。

未来，我们计划改进exchange2的性能，并检查Flang在SPEC CPU V8上的表现。其他计划包括：测试不使用LTO的编译器标志、针对Nerverse CPU进行测试，以及检查Flang在启用OpenMP时使用SPEC Speed基准测试的表现。

总结

本节课中，我们一起学习了如何分析和优化Flang编译器的性能。我们通过两个具体案例——多维数组初始化和字符串比较——探讨了识别性能瓶颈、在HLFIR层面实施优化以及评估优化效果的全过程。尽管并非所有优化都能直接带来显著的端到端性能提升，但它们在特定场景下减少了指令数量，并为整体性能改进做出了贡献。Flang的性能正在通过社区的努力持续提升。

025：VPlan的新特性

在本节课中，我们将学习过去两年里VPlan基础设施的主要改进，以及这些改进所启用的新优化。VPlan是升级和扩展循环向量化基础设施的一项持续工作。我们将回顾这些逐步引入代码库并默认启用的改进。

概述

VPlan是一个用于描述向量化候选方案的显式模型。在向量化管道的开始，我们从一个标量循环（以LLVM IR形式输入）出发，逐步构建一个初始VPlan。这个初始VPlan作为后续优化的起点。在创建初始计划的过程中，我们将抽象层次从输入的LLVM IR提升。然后，我们创建多个反映不同向量化策略的VPlan。例如，为不同的向量化因子创建不同的VPlan。一旦我们从候选集中选出最佳的VPlan，我们便逐步将其降级为向量化的LLVM IR。

回顾路线图

上一节我们介绍了VPlan的基本概念，本节中我们来看看其发展路线图。在2024年的演讲中，我们提出了五个不同的基础设施演进方向：

  1. 重构初始VPlan创建过程：将流程从“LLVM IR -> VPlan”转变为一组模块化转换。
  2. 简化VPlan执行：即从VPlan回到LLVM IR的过程。
  3. 教会VPlan计算成本，并使用这些成本直接在VPlan上做出基于成本的决策，而不再依赖传统的成本模型。
  4. 融合内联和外联向量化管道。
  5. 将执行尾部循环向量化的方式也完全转换为基于VPlan。

初始VPlan创建的改进

现在，让我们首先关注初始VPlan创建过程的改进，即如何将LLVM IR中的初始输入循环转换为初始VPlan，作为向量化和进一步优化的起点。我们将这个初始VPlan称为VPlan Zero。

以下是创建VPlan Zero的步骤：

  1. 创建初始骨架：首先创建一个不包含任何配方的平面VPlan，它只是包装了输入。循环的入口和出口块被包装在我们引入的新类型块——VPBasicBlock中。循环内的块和指令则被包装在VPBasicBlock和标量VPInstruction中。
  2. 显式建模生命周期值：离开循环的值（称为生命周期值）通过将出口块中的循环闭合SSA phi节点包装在VPRedphiRecipe中来显式建模，这些配方将离开循环的VPlan值作为显式操作数。
  3. 规范化循环：添加一个经济归纳变量配方，该变量从0步进到向量化行程计数。同时更新控制循环的退出条件，使其现在由经济归纳变量控制，确保我们要优化的循环始终是可计数的形式。
  4. 完成初始骨架：添加一些额外的块，用于在优化过程中放置各种代码。例如，引入一个新的向量化前导块来具体化向量循环中需要的一些值；添加一个中间块来调度决定在执行向量循环后是恢复执行标量循环还是跳转到原始出口块。
  5. 转换为区域块：将计划中的循环转换为专用的VPRegionBlock，其中包含我们要向量化的循环的所有块。这提升了我们工作的抽象层次，简化了后续必须处理的CFG转换，并允许我们的转换轻松识别要优化的循环。

这些是构建初始VPlan方式的一些较大变化。这让你对我们在该领域的总体思路和构建整个管道的方向有了很好的了解。还有一些标量转换我们今天没有时间详细讨论，例如线性化控制流和执行谓词化。

VPlan范围的扩展

接下来，我们看看VPlan模型范围的演变。在幻灯片上，你可以看到大约两年前VPlan的范围：它包含了入口块的部分（主要用于首次头扩展）、向量化前导块和中间块的部分，当然还有循环区域。当时，生命周期值仍然是单独建模的，这基本上限制了它只能处理具有单个出口块的循环。

如果我们看看过去两年的变化，可以看到VPlan的范围已经扩展到覆盖整个向量化骨架。我们还没有向量化尾部循环，所以处理尾部循环向量化时，这仍然是单独处理的。幻灯片上绿色高亮的部分是自2023年以来新增的。

现在，所有入口和出口块都被包装在之前提到的VPRBasicBlock中。标量循环体和标量前导块也以类似的方式建模。这允许我们对入口/出口块之间的完整控制流以及相应的运行时检查进行建模。标量前导块中的恢复值现在也使用VPRedphiRecipe显式建模，出口值则通过包装出口块的VPRBasicBlock中的VPRedphiRecipe类似地建模。这允许在VPlan执行和VPlan转换期间正确更新这些最终节点的传入值。

这种方法的一个主要好处是，它允许移除和简化各种传统代码路径中的代码。例如，这允许从向量化器类中移除处理前导对象和相应块的遗留代码，从而简化了实现。它还允许将剩余的尾部循环特定代码移动到尾部循环特定类中，使代码更清晰。

现在整个骨架都在VPlan中建模，我们可以免费获得诸如平凡运行时检查、分支和块移除等简化操作，只需在完整计划上运行现有的VPlan转换，而不仅仅是循环区域。例如，如果入口块中的最小迭代检查已知总是执行向量循环，我们可以简单地移除指向标量前导块的分支。同样，如果我们简化了中间块的条件，我们也可以移除这个平凡分支。最后，在简化之后，如果我们知道标量循环永远不会执行，我们可以通过移除包装它的VPRBasicBlock来完全移除它。

这产生了更简单、更易于审查的IR。我认为当补丁被合并时，它从各种测试中移除了大约10,000行检查行，使它们更易于审查，并且为下游传递或后续向量化留下了更少的代码来处理。

将范围扩展到包括出口块，并将构建过程分解为这组转换，也使得以一种与框架良好契合的方式支持具有提前退出的循环成为可能。只需引入另一个执行具有多个出口的循环规范化的标量转换，我们将它们再次转换为一种形式，其中向量循环由单个退出条件控制，并带有一些额外的块，以便在离开向量循环后调度到正确的出口块。出口值的显式建模意味着它们可以根据转换的需要直接调整，确保在通过提前退出离开循环时计算正确的值。

VPlan执行的简化

上一节我们介绍了VPlan范围的扩展，本节中我们来看看VPlan执行的简化，即将其降级回LLVM IR的过程。

与VPlan构建类似，这种降级可以通过将其分解为单独的转换来简化，这些转换逐步降低抽象层次。在选择最佳VPlan后，我们逐步将我们做出的隐式决策显式化，并最终转换回一个平面CFG VPlan，它更直接地映射到最终的LLVM IR。

以下是两个转换示例：

  1. 执行显式交错：除了向量化循环外，在优化过程中交错多个向量迭代通常是有利的。VPlan对一系列向量化因子和任何可能的交错因子都有效。在选择最具效益的交错计数后，降级过程中的第一步就是将交错计数应用到VPlan。对于大多数配方，这仅仅意味着将配方克隆交错计数次，并将操作数重新映射到对应交错部分的值。需要注意的一点是，这种交错保留了原始向量循环中依赖关系的顺序。类似地，我们发现后续的配方，如向量指针配方，在交错时需要一些特殊处理，最显著的是那些具有跨迭代依赖关系的配方，如归纳变量和归约。这些也由转换处理。将其作为转换执行有多个好处：一方面，它简化了所有配方的代码生成实现，这些实现以前必须负责为所有交错部分生成值；另一方面，它简化了在VPlan转换状态期间从VPlan值到相应生成的LLVM IR值的映射，变为简单的一对一映射，而不是必须处理一对多映射。更显式的实现还支持仅在应用此交错后才可能进行的额外简化。例如，在应用显式交错后，计算第一部分的向量指针变得冗余。
  2. 替换区域表示：第二个转换是将我们从具有显式区域表示的形式，转换回显式控制流。这降低了抽象层次，并将计划带回一种直接映射到LLVM IR控制流的形式，并消除了在代码生成期间处理循环区域的需要。它还支持在层次表示中进行一些额外的简化。这些循环区域要求始终提供一个经济归纳变量，并且要求是可计数的，这限制了我们在该阶段可以执行的一些转换。替换这样的区域允许我们移除不再需要的经济归纳变量，并调整循环的控制方式。例如，在RISC-V上，控制循环的条件可以替换为基于显式向量长度的条件。

向基于VPlan的成本模型过渡

另一个重大变化是向基于VPlan的成本模型过渡，即直接在VPlan上计算配方和块的成本，并使用这个基于VPlan的成本模型来做出基于成本的决策，例如选择最具效益的计划。

在接口方面，配方和块都已扩展为直接计算成本。一些配方尚未完全使用这个基于VPlan的成本模型计算成本，而是必须回退到传统成本模型。但这个过渡应该很快就会完成。

基于VPlan的成本模型可以更准确地计算成本，特别是在存在基于VPlan的转换时，以及对于一些复杂或更复杂的配方。过去几年添加了许多新配方，例如用于建模直方图归约模式的VPHistogramRecipe，这些配方仅在基于VPlan的成本模型中受支持，因为在那里支持它们非常容易，而将它们添加到传统成本模型中将非常具有挑战性。

现在，让我们也看一个转换方面的例子，这只是计划中的一个简单简化。在幻灯片上，我们有一个冗余的乘法，它将在VPlan优化期间被移除。基于VPlan的成本模型会自动处理像这样的简化，因为一旦你移除了冗余配方，在成本计算中就无需处理它们，它们已经被移除了。另一方面，要在传统成本模型中支持像这样的简化，需要大量的特殊情况处理，因为所有可能的转换或简化都需要在我们可以执行它们之前的某个时间点被预见到。在幻灯片上，你可以看到在传统成本模型中支持这种乘以1的情况所需的逻辑。这是不可扩展的，只有少数一些以前被认为重要的情况在传统成本模型中得到了处理。但正如我所说，使用基于VPlan的成本模型，我们基本上可以免费获得对任何转换的支持。

循环向量化优化管道的演变

现在，让我们通过回顾过去两年循环向量化优化管道的演变，将所有变化整合起来。

在幻灯片上，我们可以看到2023年的管道。当时的主要阶段是合法化、通用优化和特定优化。那时，所有基于成本的决策仍然由传统模型驱动。

现在我们引入了这个新的标量转换阶段，它执行我们在演讲中已经提到的初始VPlan创建。合法化和通用优化阶段也有一些新增内容，主要是为了启用新的优化，例如支持缩小交错组和在没有快速数学标志的情况下向量化浮点最小/最大归约。最后一个阶段已经泛化，以执行我们讨论过的渐进式降级，并执行向量化因子和展开因子特定的优化。

总结与进展

本节课中，我们一起学习了VPlan基础设施在过去几年的主要改进。正如我提到的，过去几年有大量的改进，这得益于贡献者数量的增长。在幻灯片上，我们可以看到过去几年每年的提交数量和独立贡献者数量。我们可以看到，无论是提交数量还是贡献者数量，都呈现出健康的增长。我认为这与VPlan基础设施的成熟相吻合，并且有效地促进了在向量化器不同部分工作的人员之间更有效的协作。

此外，过去三年中，我们能够向量化的新型循环也取得了一些进展：

  • 我们现在可以在没有快速数学的情况下向量化浮点最小/最大归约。
  • 我们初步支持向量化具有提前退出的循环。
  • 我们支持向量化针对给定谓词查找最后一个或第一个归纳变量的循环。
  • 我们现在支持所谓的部分归约，主要用于利用特定硬件指令，如AArch64上的点积指令。
  • 尾部折叠已在RISC-V上默认启用。
  • 还有许多更小的优化已经落地。

所有这些都归功于许多在向量化器上工作的贡献者，感谢所有的贡献，并希望明年能有更多进展。

本节课到此结束。谢谢。

026：更快速、更简便的方言转换驱动器

概述 📋

在本教程中，我们将介绍MLIR方言转换驱动器的重大更新：无模式回滚的“一次性”方言转换驱动器。这个新驱动器已经合并到MLIR上游代码库中。我们将探讨其核心概念、优势、API变化、性能提升以及迁移指南，旨在帮助初学者理解并开始使用这一新功能。

无回滚转换：核心概念与动机 🎯

上一节我们概述了本次更新的内容，本节中我们来深入理解其核心概念。

传统的方言转换驱动器支持“模式回滚”。这意味着，当一个模式尝试进行IR（中间表示）转换但最终失败时，驱动器可以撤销该模式所做的所有更改。然而，这种机制带来了显著的复杂性和性能开销。

新的“一次性”方言转换驱动器不允许模式回滚。所有IR更改都会立即生效。这一改变虽然对API产生了一些影响，但带来了多方面的优势：

  • 性能大幅提升：由于移除了大量用于跟踪回滚状态的数据结构，编译时间显著缩短，内存使用量大幅降低。
  • 调试更简便：每次模式应用后，IR都处于有效状态，可以轻松地通过调试或日志输出进行查看，无需处理隐藏在C++内部的状态。
  • 支持新特性：作为副作用，它更好地支持了监听器（listeners）等特性，并且使得实现“上下文感知类型转换器”等新功能成为可能，这在有回滚机制时非常困难。

动机公式可以概括为：新驱动器 = 传统驱动器 - 回滚状态管理 + 即时变更。

时间线与项目状态 📅

这个项目是一个长期工程，包含了许多前置的清理工作。目前，一次性方言转换功能已经合并。在此之前，诸如“1对N”模式支持等功能也已合并。团队还在努力使“内存到LLVM”的转换也能在此新驱动器上工作。

API差异与用户指南 📖

上一节我们了解了新驱动器的优势，本节中我们来看看用户在使用时需要注意的主要API变化。

首先，理解“有回滚”和“无回滚”的区别至关重要。

模式应用逻辑的改变

在旧驱动器中，一个模式可以：

  1. 匹配IR。
  2. 进行一些修改。
  3. 如果发现修改不适用于当前IR，则从matchAndRewrite函数返回failure()。
  4. 驱动器会完全撤销该模式所做的所有更改。

这种机制允许驱动器依次尝试多个模式。例如，驱动器可能先尝试将操作A转换为非法操作B的模式，失败回滚后，再尝试将其转换为非法操作C的模式。

在新的无回滚驱动器中，上述机制不再有效。它要求模式必须具有清晰、干净的匹配和重写阶段。如果模式匹配失败（返回failure），则绝对不能对IR进行任何修改。否则，转换可能会失败，因为后续模式无法在已被部分修改的IR上应用。

IR修改的即时性

最大的变化之一是IR修改操作现在是即时的，而非延迟的。

在旧的回滚驱动器中，诸如eraseOp（擦除操作）或replaceOp（替换操作）等方法并不会立即反映在IR中，而是将更改记录在转换驱动器内部的隐藏状态中。这样做是为了避免昂贵的IR克隆操作。

在新的无回滚驱动器中，其行为更符合用户从规范化（canonicalization）中熟悉的PatternRewriter的预期：所有更改都会立即生效。这为调试带来了巨大便利。

以下是API行为对比的摘要：

操作 回滚驱动器（旧） 无回滚驱动器（新）
replaceOp / eraseOp 延迟生效 立即生效
replaceAllUsesWith 延迟生效，有特殊语义 立即生效，语义更直观
遍历IR 可能不安全（存在未链接的临时结构） 安全（IR始终一致）

具体迁移示例与注意事项

由于修改是立即的，一些在旧驱动器中可行的代码模式在新驱动器中会导致错误。

问题1：操作使用后释放

// 旧驱动器（可行）
rewriter.replaceOp(op, newValue); // 延迟擦除
auto attr = op->getAttr("foo");   // 仍可访问‘op’
// 新驱动器（错误：use-after-free）
rewriter.replaceOp(op, newValue); // 立即擦除
auto attr = op->getAttr("foo");   // 错误！‘op’已被删除

修复方法：在修改前保存所需信息，或调整执行顺序。

// 修复：先获取属性，再替换操作
auto attr = op->getAttr("foo");
rewriter.replaceOp(op, newValue);

问题2：replaceAllUsesWith的语义
在旧驱动器中，replaceAllUsesWith是延迟执行的，并且有一个特殊行为：如果在调用它之后又创建了新的使用（use），这些新使用在最终替换时也会被更新。这容易令人困惑。

在新驱动器中，replaceAllUsesWith立即生效，行为更直观。目前，新驱动器仅支持替换所有使用场景，尚不支持replaceAllUsesExcept或replaceUsesWithIf等选择性替换的变体，但计划在未来添加。

调试与检查

为了帮助迁移，MLIR提供了一个CMake选项：-DMLIR_ENABLE_EXPENSIVE_PATTERN_API_CHECKS=ON。启用后，它会在运行时执行更严格的检查，确保模式遵守“失败时不修改IR”等关键约定，能有效捕获不兼容的代码模式。

性能对比分析 ⚡

我们介绍了API的变化，现在让我们看看这些变化带来的实际性能提升，这是引入新驱动器的首要原因。

以下基准测试来自MLIR稀疏编译器测试套件中的一个大型集成测试。该测试进行大量的稀疏张量转换和打印操作。

编译时间对比：

  • 随着问题规模（例如打印操作数量）增大，新旧驱动器的编译时间都会增加。
  • 无回滚的新驱动器始终更快。例如，当问题规模扩大80倍时，新驱动器的速度比旧驱动器快约50%。
  • 性能提升主要源于消除了大量IR重写对象的分配和维护开销。

内存使用对比：

  • 在旧驱动器中，由于IR修改被延迟且相关内存直到最后才释放，该测试峰值内存使用高达20 GB。
  • 在新驱动器中，内存被立即释放，峰值内存使用降至100 MB左右。
  • 内存使用量的巨大差异是因为新驱动器不需要维护那些用于回滚的IR重写对象栈。

性能剖析（Flame Graph）：
在旧驱动器的剖析图中，可以看到大量时间花费在分配IRRewrite对象上（例如在scf到控制流转换中分裂块时）。而在新驱动器的剖析图中，这部分开销基本消失，更多时间花在符号表查找等与回滚无关的操作上。

迁移指南与调试技巧 🛠️

了解了性能优势后，本节我们探讨如何将现有代码迁移到新驱动器，并利用其改进的调试功能。

不兼容的用例

大多数模式可以平滑迁移，但存在一些不兼容的用例。主要问题源于新驱动器只能看到最新的IR状态，而旧驱动器可以同时看到旧IR和新IR的“并排”视图。

示例：需要查看历史IR信息的模式
有一个测试用例，其模式需要根据转换前的内存操作（memref）的维度信息来计算新的索引。在旧驱动器中，即使插入了一个中间转换层（如unrealized_conversion_cast），模式仍能回溯到原始的内存操作来获取元数据。

在新驱动器中，模式只能看到最新的IR，即一个以unrealized_conversion_cast为操作数的memref.extract_strided_metadata操作，无法进一步获取原始维度信息，导致转换无法完成。

错误信息：
启用新驱动器后，如果遇到此类不兼容模式，你会看到类似错误：

convert.vector_load produced IR that cannot be legalized
...
- `memref.extract_strided_metadata` : illegal op

这清晰地指出了无法合法化的操作。

解决方案：这类情况通常需要重写整个转换流程，而不是简单的适配。有时，这类转换可能更适合使用贪婪重写驱动器（Greedy Pattern Rewrite Driver）而非方言转换驱动器。

改进的调试体验

新驱动器在调试方面有显著优势。

旧驱动器调试输出：
在转换后打印IR，你可能会看到：

  • 新旧操作并存：因为旧操作只是被标记为删除，并未立即擦除。
  • 未知的SSA值：由于块签名转换等操作是部分延迟的，打印时可能引用已分离（detached）的旧值，导致显示为<?>。

新驱动器调试输出：
在转换后打印IR，你看到的是：

  • 完全物化后的最新IR：所有操作都已就位，状态一致。
  • 清晰的转换痕迹：例如，可以清楚地看到指针如何通过unrealized_conversion_cast转换回memref类型，并与内存操作连接。

这使得通过-debug选项逐步跟踪转换过程变得非常直观和可靠。

项目经验与未来工作 🚀

本节我们总结一下这个大型重构项目的经验教训，并展望未来的发展方向。

经验教训

这个历时一年半的大型重构项目带来了一些宝贵的经验：

  1. 发送小而精的合并请求（MR）：小的MR更容易被评审和合并。
  2. 多发送无功能更改（NFC）的清理：NFC更改评审阻力小，能为后续实质性更改铺平道路。
  3. 给评审和下游项目留出时间：在发送大型MR之间留出间隔，避免评审者负担过重，并留出时间让下游项目（如Google、AMD等拥有大型代码库的团队）进行测试和反馈，这能有效发现测试用例覆盖不到的问题。
  4. 积极寻找内部用户：联系那些在内部使用相关组件的团队，让他们提前试用你的代码，可以获得宝贵的验证。

现状与呼吁

目前，方言转换驱动器可以运行在两种模式下：回滚模式和无回滚模式。
我们强烈建议大家尝试并尽可能使用无回滚模式，因为它：

  • 更快、更易用、内存占用更低。
  • 提供更好的调试信息。
  • 更好地支持上下文感知类型转换器（减少了未知SSA值的问题）。
  • 提供更准确的监听器通知（例如，操作插入通知与模式开始/结束通知正确交错）。

最终目标是弃用并移除旧的回滚驱动器。使用的人越少，这一目标就能越早实现，预计可以减少驱动器中多达50%的代码量。

未来工作

未来可能的工作方向包括：

  • 池化unrealized_conversion_cast操作：以进一步提升性能（效果待评估）。
  • 进一步收紧API：使其更贴近PatternRewriter的API。例如，当前即使一个操作仍有使用（uses），也可以擦除它，然后由驱动器插入源材料化（source materialization）。未来可能禁止此类行为，使API更严格、更安全。
  • 探索迭代顺序控制：允许用户控制区域内操作的应用顺序（例如，自底向上或先进入区域）。

总结 📝

本节课中我们一起学习了MLIR中全新的“一次性”方言转换驱动器。

我们从其核心概念——禁止模式回滚、即时变更IR——出发，理解了其带来的性能提升（更快的编译速度、更低的内存占用）和开发体验改善（更简便的调试）。我们详细分析了新旧驱动器之间的API差异，特别是IR修改的即时性带来的代码适配要求。通过迁移指南和调试对比，我们掌握了如何识别不兼容的用例并利用新工具。最后，我们分享了项目重构的经验，并展望了未来的优化方向。

现在，你可以通过在上游转换配置中设置allowPatternRollback = false来尝试这一新功能。希望本教程能帮助你顺利过渡，并享受新驱动器带来的诸多好处。

027：LLVM顾问工具介绍

在本节课中，我们将要学习一个名为LLVM顾问的统一可视化工具，它用于处理编译器生成的各类数据。我们将了解它如何简化编译数据的收集、分析和可视化过程。

概述

大家好，我是来自劳伦斯利弗莫尔国家实验室的Kevin Sala。我将介绍LLVM顾问，这是一个用于编译器产物的统一可视化工具。该项目实际上是由Miguel Carins主导的GO语言项目，我是其导师之一。

问题背景

LLVM编译器能够报告大量可用的编译数据，例如优化备注、性能剖析信息、时间统计以及诊断信息。然而，问题在于我们需要使用多个不同的编译标志来启用这些数据的报告功能。这导致输出文件格式各异，使得查找特定代码区域（例如，当我们尝试优化某些特定区域的性能时）的相关信息变得相当困难。

解决方案：LLVM顾问

为了解决上述问题，我们引入了LLVM顾问。它本质上是一个统一的基础设施工具。该工具首先通过一个名为llvm-advisor的编译器包装器，收集关于您应用程序编译过程的信息。接着，它会组织这些编译数据，将它们存储在一个文件夹中，然后进行分析和关联。最终，您可以通过一个现代化的Web界面来可视化所有这些数据。

工作原理

以下是其基本工作流程：

  1. 您拥有一个使用clang的编译或链接命令。
  2. 您只需将llvm-advisor命令放在您的编译命令之前。
  3. 该工具将收集常规编译过程中获得的所有信息。
  4. 它会在内部创建一个文件夹，存放所有信息，例如每个翻译单元的LLVM IR、汇编代码、AST和诊断信息。

可视化界面功能

最后，您可以在一个同样名为LLVM顾问的Web界面中，以统一的方式查看所有组织好的数据。

主面板（仪表盘）

界面的一侧是主面板，即仪表盘。它展示了您应用程序编译过程的一些总体信息。

以下是仪表盘显示的关键信息：

  • 您拥有的文件总数。
  • 优化备注的分布情况。
  • 诊断信息的数量及每种类型的数量。
  • 编译所花费的时间。
  • 二进制文件大小的细分。
  • 已报告的优化过程和优化备注的摘要。

性能面板

上一节我们介绍了仪表盘，本节中我们来看看性能面板。目前，该面板展示两种不同的追踪信息。

以下是性能面板包含的追踪类型：

  • 编译时间追踪：显示每个编译过程所花费的时间。
  • 运行时追踪：针对上传到GPU的应用程序，显示上传部分的运行时追踪，并展示GPU上每个操作所花费的时间。

探索面板

接下来是第三个面板，我们称之为探索面板。在这个面板中，您可以显示您应用程序的源代码（支持C++、Fortran等）。您可以在那些小标签页中启用诊断信息或优化备注的显示。它们会直接显示在报告它们的代码行上方，方便您查看。

在界面的另一侧，您可以显示针对该源文件生成的不同输出。

以下是可显示的生成输出类型：

  • LLVM IR：可以显示LLVM中间表示。
  • 汇编代码：可以显示生成的汇编代码。

适用人群

那么，谁能从LLVM顾问中受益呢？

我们认为以下人群可以从中获益：

  • LLVM高级开发者：可以利用它减少翻阅优化备注文件、寻找感兴趣特定部分所花费的时间。
  • LLVM新手：可以更容易地了解LLVM基础设施和编译过程。
  • 教育工作者和大学教授：可以将其用于教学，帮助学生理解编译过程。

总结与资源

本节课中，我们一起学习了LLVM顾问工具，它如何通过统一的Web界面来收集、组织和可视化复杂的编译数据，从而帮助不同背景的开发者更高效地进行工作。

如果您有任何问题，欢迎在明天的海报展示环节来访，我将在那里进行展示。正如所说，这是由米尼奥大学的Miguel Cardis学生主导的项目，您也可以联系他。最后，我在这里留下了二维码，如果您想尝试LLVM顾问的实时预览，或者查看正在审核中的PR，都可以扫描获取。任何评论或反馈都非常欢迎。非常感谢。

028：循环优化与未来展望

概述

在本节课中，我们将学习LLVM中循环优化的现状、挑战以及未来的发展方向。我们将探讨如何通过改进依赖分析和类型信息提取来提升循环优化的性能，并了解社区在推动这些优化成为默认设置过程中所面临的挑战。

循环优化的背景与重要性

上一节我们介绍了循环优化在提升程序性能中的关键作用。本节中，我们来看看为什么LLVM需要加强这方面的能力。

许多专有编译器已经实现了循环分布、循环交换等优化，这使得LLVM在某些基准测试中表现落后。例如，在TSVC基准测试中，由于缺少循环分布和交换优化，LLVM的性能可能落后专有编译器高达35倍。

TSVC虽然是一个人工合成的基准测试，但其揭示的模式在真实工作负载中同样存在。如果编译器在TSVC中错过了某些优化模式，那么在实际应用中也可能会错过。

循环优化的实际影响

以下是我们在真实工作负载中观察到的循环优化带来的性能提升示例：

  • Spec CPU 2017/2006/2000：多个基准测试显示性能有显著提升。
  • Livermore Loops：经典的科学计算内核，优化效果明显。
  • Polybench/RAJA/Perf/Geekbench：涵盖多种计算模式，验证了优化的广泛适用性。

需要说明的是，这些数据基于特定的硬件平台（AX 60），实际效果可能因环境而异。此外，最大的性能提升往往来自于多种变换的组合应用，例如循环分布与循环交换的结合。

当前挑战：“不可能的任务”

我们的目标是在去年的LLVM开发者大会后，让循环交换优化在默认优化管道（如-O2及以上）中启用。然而，由于大量的错误修复、技术讨论和固有挑战，我们未能如期实现这个目标。

目前，LLVM中许多经典的循环优化（如循环分布、融合、交换）虽然存在，但默认是关闭的。这造成了巨大的技术债务。我们的策略是首先在Flang（Fortran前端）中启用循环交换，因为科学计算Fortran代码对这些优化非常敏感，能提供良好的测试用例。之后，我们计划在Clang中做同样的事情，并着手研究下一个目标：循环分布优化。

技术难点：依赖分析与线性化

循环交换等优化需要确保变换后的程序语义正确，这严重依赖依赖分析。依赖分析的核心是判断不同的数组访问（读/写）是否指向相同的内存位置。

为了高效地进行依赖测试，我们需要识别出数组访问的下标。给定扁平的LLVM IR，我们通过线性化过程来恢复这些下标信息。例如，对于源代码中的访问 A[i][j]，线性化分析需要从一系列 getelementptr 指令中解析出下标 i 和 j。

当前面临的主要问题包括：

  1. IR简化：LLVM的IR简化过程可能会移除 getelementptr 指令的类型信息，而这正是线性化分析所依赖的。
  2. 类型信息提取：我们需要从全局变量和动态分配数组中提取维度信息，并将其表达在IR中，以辅助优化器做出更好的决策。

对于静态声明的全局数组，我们可以通过附加元数据来提供维度信息，例如：

!0 = !{!"array_info", ptr @A, i64 0, i64 42} ; 表示数组A的索引范围是[0, 42)

对于动态数组（包括Fortran的假定形状数组），我们可能需要插入额外的assume指令来向优化器传递边界约束，例如“索引 i 必须小于42”。

社区与开发流程的启示

在推进此项工作的过程中，我们也获得了一些关于LLVM开发流程的启示。

LLVM社区已经变得更加成熟。一项重要的政策变化是：当添加一个新的编译通道时，目标应是尽快将其纳入默认优化管道，并持续进行增量开发，且该通道不应存在已知的正确性问题。

然而，历史遗留问题是，许多循环优化通道因未能遵循此流程而默认关闭。这导致我们现在需要花费大量精力去“解耦”这些通道中堆积的功能，修复性能与正确性问题，这并非最高效的方式。

此外，编译时间始终是社区关注的焦点。由于这些优化默认关闭，其编译时间开销未被有效监控，导致问题不断累积。我们需要对所有感兴趣的循环优化进行同样的处理：测量、改进编译时间，并修复正确性问题。

未来计划与总结

本节课中我们一起学习了LLVM循环优化的现状与核心挑战。我们的未来计划清晰分为几步：

  1. 修复依赖分析：这是当前启用循环交换的主要障碍。我们需要解决一系列相关问题，特别是处理整数溢出等边界情况。
  2. 启用循环交换：在依赖分析稳固后，在Flang及随后在Clang中默认启用循环交换优化。
  3. 增强信息表达：通过从源码提取数组维度信息并编码到IR中，为优化器提供更丰富的上下文，以解锁更大的性能潜力。
  4. 建设社区：循环优化需要更广泛的社区参与。我们鼓励大家参与每月举行的“循环优化工作组”会议，共同review代码，推动进展。

关于MLIR与LLVM IR的路线，我们认为基于MLIR的循环优化仍然较远，涉及大量基础设施迁移和成本模型重建。因此，在可预见的未来，继续在LLVM IR层面启用和完善这些优化是务实且必要的选择。

总结：提升LLVM的循环优化能力是一项持续的任务，需要解决底层的技术难题（如依赖分析），改进开发流程，并依靠社区的力量共同推进。我们相信，通过逐步启用和优化这些变换，LLVM能够在高性能计算领域更具竞争力。

029：利用IR2Vec嵌入增强MLGO内联

在本节课中，我们将学习如何利用IR2Vec嵌入技术来增强LLVM中基于机器学习的优化（MLGO），特别是函数内联优化。我们将探讨如何更好地表示程序以供机器学习应用使用，并详细介绍IR2Vec的工作原理及其在LLVM中的集成。

概述

我们试图回答的核心问题是：如何为机器学习应用更好地表示程序？为此，我们有两种主要方法：基于特征工程的方法和基于嵌入的方法。特征工程依赖于人类专家确定对特定优化相关的特征，而嵌入则是一种学习到的表示，旨在自动发现重要信息，从而以任务无关的方式驱动优化决策。

嵌入技术简介

嵌入是一种学习到的表示。本质上，我们使用一个机器学习模型，将对象作为输入，并将其投影到一个N维空间中的浮点向量，这些向量捕获了有意义的信息。在自然语言处理中，嵌入可以捕获诸如“男人”与“国王”之间的关系类似于“女人”与“王后”之间的关系。我们的目标是：能否为LLVM IR做类似的事情？我们能否完全消除特征工程？

为此，我们提出了一种称为IR2Vec的方法。这是一种基于LLVM IR的方法，我们利用程序分析信息来驱动嵌入的生成。顾名思义，这是一种分布式编码，其中语义含义分布在向量的各个组成部分中，而不是典型的基于特征的表示。该方法采用自底向上的方式：我们从学习实体级别的表示开始，然后聚合它们，最终得到IR不同区域的表示。

IR2Vec工作原理

让我们简要概述这个过程。考虑一个典型的LLVM IR。

该IR中的每条指令 I 被分解为一组三元组。每个三元组的形式为 (H, R, T)，其中 H 代表头，T 代表尾，R 代表连接头与尾的关系。

我们主要定义了三种关系：

  1. 类型关系：将操作码与指令的类型关联起来。
  2. 下一条指令关系：将当前指令的操作码与下一条指令的操作码关联起来。
  3. 参数关系：将操作码与指令的参数关联起来。

这个过程在大量的IR语料库中对每条指令进行。我们使用一个简单的神经网络进行训练。

这里我们使用了一个称为TransE的翻译表示学习模型。该模型学习的关系是：头 H 加上关系 R 应近似等于尾 T，即 H + R ≈ T。

训练过程结束后，我们为训练期间遇到的所有实体获得了一个表示。所有这些被整合在一起，形成了一个我们称之为种子嵌入词汇表的东西。这个词汇表的训练是离线一次性完成的。一旦有了这个词汇表，我们就能够推导出不同级别的IR表示。

你可以将这个词汇表想象成一个简单的JSON文件，包含大约100个实体，其中键是实体的名称，值是该实体的向量表示。

生成指令表示

一旦有了词汇表，我们通过查表获取操作码、类型和参数的表示，并将它们组合起来得到一条指令的表示。

这个过程结束后，我们能够为IR中的每条指令生成表示。这种生成表示的方式我们称之为符号编码。

但我们并不止步于此，我们还希望编码一些流信息。

正如之前所见，指令中的操作数会获得一个通用的表示（如指针、变量、常量等）。但我们希望将其特化，并在此处编码上下文的概念。我们通过使用流信息来实现这一点。

假设我们像这样推导出指令 I2 的表示。对于指令 I4，我们考虑所有可能到达该点的定义，并利用这些表示的组合来推导 I4 的表示。如图所示，来自 I2、I3 和 I9 的定义都可能到达这里。我们保守地使用所有三个定义来推导 I4 的表示。这个过程可以对所有指令重复进行。

我们可以最终编码更多从定义-使用分析中推导出的上下文信息来获得表示。这就是我们所说的流感知编码。

由此可见，这种生成表示的方式更简单，同时也更具表现力。例如，你可以考虑编码整个函数的表示，并用该表示来表示调用点，这反过来可以模拟内联的效果。如果你知道某个特定定义到达一个使用点的可能性，你还可以进一步特化，编码这种可能性并相应地加权各个组成部分，这通常可以通过使用性能剖析信息来推导。

聚合区域表示

一旦我们有了指令级别的表示，就可以使用不同的聚合器来组合指令的表示，以获得代码中不同区域的表示，例如函数、循环、基本块等。

你可以考虑各种聚合器，比如简单的线性组合、求平均，或者训练模型来学习更好的聚合方式以获得更好的表示。

这本质上就是我们基于LLVM IR推导出的IR2Vec。它也有衍生版本，例如我们为LLVM的机器IR（主要考虑后端应用）推导的表示，称为MIR2Vec。我们还有另一个针对二进制文件的变体。如果你有二进制文件并想在二进制级别进行优化或转换，可以使用从VEX IR（Valgrind的IR）推导出的类似表示。

LLVM中的IR2Vec

现在，让我们谈谈LLVM中IR2Vec的现状。目前，IR2Vec和MIR2Vec已在LLVM上游代码库中可用。我们发起了一个RFC并进行了讨论，然后开始以一系列增量补丁的形式将其上游化。

IR2Vec位于 LLVM/lib/Analysis 目录下，MIR2Vec位于 LLVM/lib/CodeGen 目录下。我们还有一个独立的工具叫 llvm-ir2vec，位于 LLVM/tools 目录下，它可以帮助生成嵌入和进行词汇表训练。这里的想法是，你可以使用这个工具或即将推出的Python库来驱动你的训练、在Python中生成嵌入以进行训练；在推理时，LLVM Pass可以直接从IR分析中消费这些嵌入。原始源代码也是开源的。

IR2Vec包含三个主要组件：

  1. 词汇表分析：读取包含种子嵌入词汇表的JSON文件，并将此信息提供给嵌入类。
  2. 嵌入类：包含如何生成表示、如何生成嵌入的核心功能。嵌入类使用词汇表分析提供的结果。
  3. 嵌入本身：只是对标准向量的包装。可以将其视为N维表示被包装为嵌入。

在LLVM Pass中如何使用：

  1. 首先运行词汇表分析Pass，获取词汇表结果。
  2. 创建一个嵌入实例。
  3. 一旦有了嵌入实例，你就能够查询并生成不同级别的表示。我们目前提供函数、基本块和指令级别的表示作为起点。

使用独立工具，你可以运行 llvm-ir2vec 并使用 embedding 子命令来生成嵌入，或者使用 triplet 子命令来生成用于训练的三元组。你可以查看LLVM工具中提供的脚本，了解如何在大规模语料库上自动生成三元组。

应用与成果

截至目前，我们已经使用IR2Vec来驱动不同的、基于机器学习的优化模型，例如循环分布、阶段排序和寄存器分配。这些都是研究项目。在LLVM中，我们有两个基于机器学习的引导优化：函数内联和寄存器驱逐决策（在贪心寄存器分配器中）。接下来我将简要介绍我们如何在函数内联中使用IR2Vec。展望未来，我们计划使用MIR2Vec来驱动贪心寄存器分配器中的驱逐决策。

在使用ML内联顾问的工作中，与使用基于特征的MLGO表示相比，我们在不同工作负载上看到了约2%到5%的性能提升。作为第一步，我们只是将嵌入与特征连接起来并训练了一个模型。该模型在一个包含约50,000个模块的内部数据中心二进制文件上训练了约2000万步，使用了PPO策略。这个训练好的模型被用于评估不同工作负载的性能，并且该模型被训练为优化代码大小。这是我们目前的技术水平。

展望未来，我们的想法是开始移除特征，并尽可能用嵌入替换它们。目前，我们有大约32个不同的特征驱动着机器学习内联决策。考虑到我们使用了嵌入，这些特征中的大多数可能是冗余的，显然可以被移除。

挑战与未来方向

接下来谈谈挑战和未来的道路。首先，正如你所想象的，像流感知嵌入这样的东西可能会遇到循环依赖。这在考虑PHI指令时非常典型，其中可能有多个定义以循环方式到达。那么如何解决呢？事实证明，当我们对流感知方程建模时，它表现为一组联立方程。一个非常明显的解决方案是使用某种线性求解器来解决它。如果我们想使用线性求解器，是应该使用像Eigen这样的现有库，还是应该编写一个简单的手写代码并将其放在LLVM工具中？或者我们可以使用精度较低但能编码流感知本质的迭代解决方案。

另一个明显的事情是编码更多信息。例如，我们可以从内存依赖、内存别名分析和内存SSA分析开始。LLVM中的这些分析甚至可以使这些流感知嵌入的构建更加优雅。但本质上，它们是非常保守的，因为它们都是为优化设计的，不能容忍假阴性。然而，嵌入本身可以承受假阴性。因此，如果我们能稍微放宽其保守性，允许一些假阴性，我们就可以编码更好的信息。如果是这样，我们应该为嵌入生成的目的专门定制这些基础设施，还是可以进行某种内存剖析，并据此决定考虑哪些依赖？

当涉及到MIR词汇表时，我们有特定于目标的指令操作码和其他操作数。这里与典型的基于LLVM IR的表示的一个主要区别是词汇表的大小。IR2Vec词汇表总共只有大约100个或更少的实体（包括操作码、类型、操作数等）。相反，在MIR（例如x86_64特定词汇表）的情况下，我们有大约7000个操作码可用。我们尝试通过分组操作码来减小这个大小。例如，你可以将像 ADD32rr 这样的操作码规范化成典型的 ADD，并最终将这些额外的前缀和后缀分开编码。问题是，x86提供了系统的方法来做到这一点，因为它使用适当的TableGen规则来定义前缀和后缀，我们可以在规范化操作码时利用这些。但这并不是通用的，不能跨不同架构。因此，在这种情况下，我们如何进行分组或规范化是另一个可能需要回答的问题。

最后，我们还计划自动化词汇表的生成。可以想象，词汇表在不断演变，编译器也在不断演变。与相对稳定的LLVM IR相比，我们在MIR中看到越来越多的操作码和操作数变化。在这种情况下，我们可能需要定期训练词汇表并使其可用。一个明显的方法是使用构建机器人。我们可以将这个过程解耦为两个阶段：

  1. 生成三元组：可以在LLVM代码库上完成。我们可以使用LLVM本身作为生成三元组的数据源。这个过程本身是轻量级的，使用大约64个CPU核心需要大约10分钟。这也可以作为集成测试。
  2. 训练模型获取词汇表：一旦有了数据，我们就可以训练模型来获得实际的词汇表。我们实际上可以以较低的频率训练模型，并使词汇表可用。

总结

本节课我们一起学习了IR2Vec嵌入，这是一种学习到的表示，已被集成到LLVM上游。我们解决了不同的性能相关问题，你可以通过分叉LLVM来使用它。其目标是减少特征工程的工作量，并改进当前的MLGO基础设施。

我们已经看到了它在研究和实际生产环境中的潜力。接下来的步骤是减少当前使用的特征，并开始使用MIR2Vec来驱动贪心寄存器分配器中的驱逐决策，以及自动化词汇表训练流程。

030：如何修改LLVM IR

在本节课中，我们将学习如何对LLVM进行比单个补丁更复杂的修改，包括添加、移除或修改某些功能。我们将探讨在此类变更中应遵循的最佳实践和流程。

概述

LLVM是一个庞大而复杂的编译器基础设施。当你需要对其进行实质性修改时，遵循一个清晰的流程至关重要。这不仅有助于确保技术方案的正确性，也能提高你的修改被开源社区接受的可能性。本节课将引导你完成从构思到部署的完整过程。

设计阶段：构思你的变更

上一节我们介绍了课程目标，本节中我们来看看如何开始设计你的变更。

首先，你不应该一开始就编写大量代码并试图直接实现它。优秀的程序员更关心数据结构及其关系，而非代码本身。

LLVM IR中的意义并非来自单个元素，而是来自它们彼此之间的关系。例如，指令之间存在直接的数据依赖关系，元数据可能编码了指针间的别名信息，调用指令则隐式地定义了其前后代码的关系。

当你向IR中添加新内容时，需要思考它将如何改变这些元素之间的关系。如果你的变更导致许多关系发生潜在变化，那么它就是侵入性的。侵入性的设计会迫使其他开发者考虑你的变更，从而增加复杂性，应尽量避免。

以下是一个侵入性设计的例子：

  • 假设你决定为某些函数添加一个名为 no_overflow_at_all 的属性，以忽略所有溢出行为。
  • 虽然IR层面的改动很小，但所有考虑溢出行为的代码站点都必须处理这个新属性。
  • 未来任何编写依赖溢出行为的新代码的开发者，也必须考虑你的属性。
  • 即使没有立即给他人带来额外工作，这仍然是一个侵入性设计。

关于数据结构，LLVM主要有两个层次结构：Value 层次结构（包含指令、常量等）和 Metadata 层次结构（包含调试信息等额外数据）。尽量让你的修改位于这两个层次结构内，这样可以免费获得序列化、哈希等支持。如果必须添加新的数据结构，请考虑计算复杂度、编译时间影响和内存使用，并优先复用LLVM ADT（抽象数据类型）目录中的现有组件。

撰写RFC：寻求社区反馈

在理论设计之后，你需要将想法转化为RFC（征求意见稿），并在社区论坛上寻求批准。

即使对于较小的变更，撰写RFC也很有必要。LLVM代码库非常庞大，无人能完全理解其所有细节。将你的提案发布在论坛上，可以让了解编译器不同角落的开发者告诉你潜在的问题，从而节省大量时间。

以下是撰写RFC时的一些要点：

  • 提供对技术变更的简洁而详细的解释。
  • 阐明变更的动机和目标。
  • 明确变更的受益者（目标受众），这有助于证明将其纳入开源编译器的合理性。
  • 尝试量化维护成本和潜在问题，以便社区进行成本效益评估。
  • 开篇段落至关重要，需要吸引读者的注意力，让他们有兴趣阅读全文。

在收到RFC反馈后，你可能会遇到几种类型：

  1. 技术性反馈：指出你方案中的细节问题、适用范围或更好的实现方式。
  2. 意愿性反馈：社区成员表达对你的想法的支持或对方向的担忧。获得支持性反馈非常重要，它表明存在感兴趣的受众。
  3. 沉默：如果没有收到反馈，这可能意味着受众较小或无人关注，但这不代表你的想法被拒绝。对于侵入性较小的独立变更，你可以继续推进；对于侵入性大的变更，这可能意味着难以证明其合理性。

如果遇到沟通困难，无法在RFC上达成一致，强烈建议参加LLVM开发者大会。面对面的交流更直接、坦诚，有助于理解他人的观点和动机。

实现阶段：构建原型与编写代码

假设你的提案已获得某种共识，并且有受众支持，接下来可以开始构建原型。

在原型阶段，目的不是编写高质量代码，而是快速验证想法、连接编译器中被修改的行为与其他部分，并尽早发现可能存在的致命问题。

以下是一些原型开发建议：

  • 使用 llvm-reduce 工具来生成最小的测试用例，这些用例在未来可作为冒烟测试、回归测试或编写上游测试的基础。
  • 在开发过程中，拥有一个独立于LLVM的质量评估目标。例如，如果你修改了底层架构，可以证明生成的二进制文件与之前完全相同，这能提供很大的信心。
  • 开启LTO（链接时优化）进行测试，因为它会让编译器的每个部分都相互摩擦，暴露出各种奇怪的行为和问题。
  • 如果你修改了文本IR，请考虑其可读性。例如，使调试记录缩进，让眼睛能轻松区分真实指令和元数据。
  • 如果更改了文本IR，也必须更改位码（bitcode）。不必害怕修改位码，其防护机制很强，会引导你做出正确的选择，但需注意大小变化和向后兼容性要求。

代码审查：准备与提交补丁

当原型完成并经过评估后，你需要准备并提交代码进行审查。

提交补丁的唯一正确方式是提交小型、增量的补丁。审查时间与变更规模呈非线性增长。一个触及编译器多个部分的大型补丁会让审阅者需要考虑的细节呈指数级增长，从而难以获得及时审查。

以下是如何拆分补丁的建议：

  • 理想情况：如果你的项目结构允许，将其拆分为一系列独立的小补丁。例如，先修改基础数据类型，然后是底层架构，接着是各个Pass的插桩，最后是前端（如Clang）的改动。
  • 大型算法：如果必须添加一个大型独立算法，可以尝试将其拆分为多个小模块，分别添加单元测试，最后再用一个“顶石”补丁将它们组合起来。
  • 多处调用点：如果你需要修改大量独立的调用点，可以预先通过机械性的补丁将它们“标准化”，这样后续实现核心功能的补丁就只是对大量相似代码做相同修改，便于审阅。
  • 无法拆分：如果确实无法拆分（例如一个完整的新优化Pass），可以寻求许可，先以禁用状态提交代码，然后逐步启用它，并随着启用范围扩大而添加测试。

请记住，社区中可用的审查时间非常有限。优化你的补丁以方便审阅者高效工作，是提高合并成功率的关键。

部署与维护：合并代码与长期责任

最后，假设你的变更已经过审查并获得批准，接下来是部署阶段。

你不应该一次性推送所有补丁。总会有边缘案例和未知问题在最后时刻出现。你需要像提交审查时一样，逐步部署你的补丁。

部署时需考虑两点：

  1. 回退的难易程度：如果你的补丁触及数百个文件，回退将非常困难。
  2. 影响范围：你的新代码或Pass是否会突然处理编译器的所有输入？能否更渐进地引入？

缓慢合并补丁（例如20个补丁在20周内合并）是理想情况，这为持续集成（CI）测试和其他开发者适应变更留出了时间。如果必须进行无法渐进部署的“阶跃式”变更（例如移除调试内部函数），请提前公告，提供简单的开关（逃逸舱口），并尽量保持差异（diff）最小化，以便于回退。

变更合并后，维护工作随之而来：

  • 为你添加的功能编写文档。
  • 如果添加的是通用设施，社区可能会在其基础上进行扩展。
  • 如果你添加的是非常特定或小众的功能，社区期望你能持续维护它。如果某项功能不断出现问题却无人维护，社区可能会因其成为维护负担而将其移除。

总结

本节课中我们一起学习了修改LLVM IR的完整流程。关键点总结如下：

  • 设计：使你的变更非侵入性、渐进演化且符合人体工学。
  • RFC：尽早寻求反馈，并以真诚的态度回应所有反馈，包括负面意见。
  • 原型：快速构建原型以发现所有潜在问题。
  • 审查：为审阅者优化你的补丁，认识到社区审查时间的稀缺性，充分准备并提交增量补丁。
  • 部署：以渐进方式提交和部署补丁。
  • 维护：承诺对你引入的变更进行长期维护，以建立信任。

遵循这一流程，可以帮助你更有效、更顺利地将有价值的改进贡献给LLVM社区。

031：MARCO - 基于MLIR的Modelica编译器

在本教程中，我们将学习MARCO项目，这是一个基于MLIR的编译器，专门用于处理名为Modelica的领域特定语言。我们将了解其设计动机、核心架构、引入的新方言以及一些关键的优化技术。

项目动机与Modelica语言简介

上一节我们概述了MARCO项目。本节中，我们来看看驱动该项目开发的核心动机以及它所针对的Modelica语言。

Modelica是一种非因果、面向对象和组件的语言，用于描述微分代数方程系统。例如，描述一个电容器可以通过定义其两个引脚，并使用“连接器”的概念将它们连接起来。随后，可以描述支配该系统的方程，在此例中即电容器的物理定律。当然，还可以向系统中添加其他组件以构建更复杂的场景。

关于Modelica代码，有两点需要注意：

  1. 代码中可能包含不常见的运算符，例如时间导数运算符。
  2. 程序主体由方程构成，而非赋值语句。这些方程之间没有强制性的执行顺序，赋值操作也不是显式的。

换句话说，Modelica是一种声明式语言，而非编程语言。因此，以下代码在语义上并不代表一个循环，而是多个方程的同时迭代（在此例中是100个）。

for i in 1:100 loop
  x[i] = y[i] + z[i];
end for;

这一特性对编译器工具至关重要。

那么，为什么我们需要构建一个新的编译器呢？事实证明，现有编译器对上述最后一个特性（数组方程和循环）的处理并不理想。它们通常的做法是展开这些循环，并将微分代数方程算法应用于展开后的源代码。这会导致其内部表示变得极其庞大。

虽然展开有时可能带来性能优势，但它对编译性能的影响是灾难性的，编译时间会迅速变得不可行（例如长达数天）。此外，还会产生巨大的二进制文件（例如GB级别）。

因此，MARCO项目的主要动机是超越当前处理代数方程系统的方法，设计能够推理数组变量和方程循环（我们简称为数组方程）的新算法。其核心目标是能够模拟通常由这种结构表示的大规模系统。

为什么选择MLIR？

上一节我们了解了MARCO要解决的问题。本节中，我们来看看为什么选择MLIR作为实现基础。

使用MLIR的优势包括：

  • 利用开源编译器基础设施。
  • 共享跨不同领域的概念和优化。
  • 与LLVM后端有更紧密的集成。
  • 可能更容易为自定义架构生成代码。
  • 更好的可调试性。

MARCO的架构与编译流程

了解了MLIR的优势后，我们来看看MARCO的整体架构和内部表示是如何构建的。

MARCO的结构对于基于LLVM的编译器来说是典型的。它有一个前端处理Modelica代码，然后进入中间表示，最后链接到构成运行时系统的库。然而，MARCO也与Clang驱动程序集成，目的是接收目标和优化特定信息，并使用Clang前端处理C代码（Modelica规范允许通过外部函数实现某些功能）。

在MARCO前端的编译流水线中，编译过程分为三个主要部分：

  1. 规范化阶段：处理一些语言特定特性，并为流水线后续部分建立某些假设。
  2. 数学处理阶段：涵盖微分代数方程处理的数学方面，例如执行因果化（许多新算法所在之处），并对系统应用数值积分。
  3. 优化与 lowering 阶段：更传统的部分，应用一些优化，并将各种方言逐步降低到LLVM IR，然后交给LLVM的中端和后端。

目前，前端总共有85个Pass，但我们不会在此详述所有。

MARCO引入的新方言

从流水线中可以看到，我们引入了一些新的方言（位于左侧）。现在，我将简要介绍它们，重点说明某些设计决策背后的原因，这些可能对社区有益。

最相关的方言无疑是 bmodelica 方言。你可能会问为什么是“b”。实际上，目前的MARCO并不直接接受完整的Modelica代码，而是接受其一个子集，称为Base Modelica（因此是“b”）。完整的Modelica代码会先由另一个编译器处理以移除面向对象的特性，然后MARCO从Base Modelica开始处理。理论上所有步骤都可以在MARCO内完成，但目前人力有限。

bmodelica方言的主要概念是表示微分代数方程系统。其核心概念是模型，我们通过一个具有名称的单一区域操作来表示。在该区域内，我们描述变量和方程。

以下是关于变量和方程表示的一些关键设计：

变量表示：我们有一个明确的设计目标，即保持与其他Pass的互操作性。在MLIR生态系统中，某些操作具有“从上方隔离”的特性，这阻止了区域内的操作引用区域外定义的SSA值。为了真正对IR中可能使用的任何其他方言开放，我们必须考虑这一点。解决方案是对模型变量采用内存语义，并提供一些操作在需要时在内存语义和SSA语义之间切换。然后，根据在IR中使用的位置，对这些操作（如变量声明和绑定）进行不同的 lowering。例如，在模型内部，你得到全局变量；如果在函数内部声明变量，你可能得到可被优化的分配操作。

方程表示：考虑到MARCO的转换过程，方程的表示也并非直截了当。我们的想法是从一开始就避免复杂化，因为某些分析的结果可能依赖于特定方程的索引（记住，方程可能基于多个索引进行迭代）。为了实现这一点，方程体与实际实例化被分开。我们称前者为方程模板，后者为方程实例。方程模板在由实例指定的索引上是参数化的。此外，我们使用内部开发的高效归档实现来支持索引范围的压缩。在这种情况下，我们更倾向于使用SSA语义，因为我们对IR结构有完全的控制权，嵌套问题不再存在。

除了bmodelica，我们还有其他几个方言，我将快速介绍：

  • eda、kinsol 和 sundials 方言：用于与外部求解器集成（例如SUNDIALS套件中的IDA和KINSOL）。它们在模型求解阶段（编译流水线的中间部分）使用，用于提供隐式积分方法或数值求解系统中仍然存在的循环。
  • runtime 方言：与MARCO的运行时环境接口。它提供一些示例函数（如正弦、余弦的实现），也用于提供编译模型所期望的运行时环境中的某些核心扩展。

优化示例：方程并行化

上一节我们介绍了MARCO的核心方言。本节中，我们通过一个具体的优化示例来展示其应用，你将看到其中两个方言的协作。

这个优化的想法是并行化独立方程。想象你有一个依赖关系图，其中方程1只有在方程0计算完成后才能计算。你可以做的是计算方程之间的依赖图，并理解例如组1内的所有方程可以以独立的方式计算。这看起来可能简单且低效，因为只有两个方程用两个线程。但请记住，这两个方程可能是数组方程，并且彼此独立。因此，你可以做的是分割所有索引范围，并将所有方程分配给多个核心。这就是主要思想，并不复杂。

从IR的角度来看，考虑一个简单例子，你有一些不透明的块，其中只包含关于依赖关系的知识（哪个方程读取/写入哪个变量，需要哪些其他变量），并且有一个额外的属性说明该块是否可以与其他块并行化。

首先，我们应用依赖图分析并包装这些组。ScheduledBlocks操作就代表了上一张幻灯片中显示的组。然后，应用一个特定的转换来查看这些块是否可以被并行化。如果可以，它会插入一些由运行时系统管理的运行时调度对象，这些对象将负责分发方程、分割索引并在多个线程间调度方程。

性能展示

现在，让我们看一个展示项目性能的例子。

考虑一个硅芯片的热模型，这是一个非常简单的模型。你可以用三个维度参数化你的芯片，根据这些参数控制体积单元的数量。底部有一个固定的电源，顶部有一个固定的冷却表面。

我们使用谱方法，测量了MARCO和另一个开源编译器OpenModelica的编译时间和模拟时间。

从结果可以看出：

  • 编译时间：MARCO在编译时间上表现出色，因为它利用了编译过程中那些数组感知算法，基本上实现了恒定的编译时间，这与OpenModelica等编译器非常不同。
  • 模拟时间：MARCO也带来了可观的收益。并且你会注意到，在某个点，运行时系统会自动激活多线程，因为它理解到在系统扩展到一定程度时，激活多核调度会变得有益。

当前状态、挑战与总结

本节课中我们一起学习了MARCO编译器的核心内容。最后，我们来总结其现状、遇到的挑战并展望未来。

当前状态：MARCO仍然是一个原型编译器。虽然已经开发了五年，但仍有工作要做。例如，我们对转向自定义架构非常感兴趣，即使只是使用GPU也将是非常好的事情，特别是对于MARCO的用例。

对社区的启示与挑战：在开发过程中遇到了一些困难，可能对社区有借鉴意义：

  1. 与Clang驱动程序的集成：并不容易，因为它是一个庞大的代码库。虽然未来可能会改善，但目前仍需要一些复制粘贴，并且关于LLVM源码修改仍有一些疑问。
  2. MLIR操作的特性（Attribute/Property）表：它们很好，但需要手动更新。流水线中的某些点在规模上曾呈二次增长（这些问题已被修复），但我在想我们是否能做得更好，例如能否通过Trait自动将特性表附加到操作上，并拥有自动更新它们的基础设施？
  3. 基于接口的成本模型：我们非常有兴趣开发一些基于接口的机制来为操作建立成本模型。因为目前自动分组并行独立方程的机制非常朴素。我们希望考虑每个方程相对于其他方程的计算成本，以便更好地在多个核心间平衡计算。但这带来了许多问题。
  4. 早期访问（Early Access）：这在开发的最初几天是个更大的问题，我知道相关工作正在进行中，但目前它被弃用了，我们现在没问题。

总结：MARCO是一个基于MLIR的、针对Modelica语言的编译器项目，旨在通过创新的数组感知算法高效处理大规模微分代数方程系统。它引入了多个专用方言，并在编译性能和并行计算方面展示了潜力。项目欢迎贡献。

032：Go编译器的LLVM后端

概述

在本教程中，我们将学习如何为Go语言构建一个基于LLVM的后端编译器。我们将探讨其背后的动机、标准Go编译器的基础知识、实现方法，特别是如何支持Go的垃圾回收机制，并最终评估其性能表现。

章节 1： 动机与背景

上一节我们介绍了本教程的主题，本节中我们来看看为什么需要为Go语言优化性能。

Go语言被设计为一种易于学习的编程语言。其编译器编译速度非常快，并能生成较小的二进制文件。然而，性能并非Go社区的首要关注点。

在阿里巴巴，Go是后端服务中使用最广泛的编程语言，每天消耗大量的CPU核心资源。为了降低成本，我们有许多团队从不同方面尝试提升Go的性能。例如，有团队致力于RPC框架，有团队优化Go标准库，还有团队研究运行时内存管理。

作为编译器团队，我们最初尝试为Go编译器添加一些优化。我们曾试图将LLVM中的一些循环优化迁移到Go编译器中，但由于其复杂性，我们放弃了。

随后，我们开始研究基于LLVM的实现方案，尝试用LLVM来编译Go代码。

章节 2： 现有方案与挑战

上一节我们介绍了性能优化的需求，本节中我们来审视现有的Go编译器工具链及其挑战。

我们首先考察了TinyGo。TinyGo是为嵌入式环境设计的，并不适用于服务器端应用程序。它拥有独立的运行时和SDK，这意味着绝大多数后端服务无法使用TinyGo构建。

我们也检查了Gollvm项目，但它存在与TinyGo相同的问题。

因此，我们开始尝试自己构建一些东西。

以下是标准Go工具链的构成：

  • 一个Go程序不仅包含Go源代码。
  • 它还包含一些使用Go特定Plan9语法的汇编代码，这类汇编代码只能由Go汇编器（go asm）处理。
  • C源代码也通过外部函数接口（FFI）在Go程序中被广泛使用。
  • 此外，Go程序还需要与Go运行时链接。Go运行时依赖于编译器生成的一些模块数据来支持其操作，这类元数据用于支持栈追踪、垃圾回收和反射API等。

回顾标准工具链，如果你计划用LLVM实现Go编译器，可能需要实现所有相关组件，例如Go ABI、Go特定的对象格式、垃圾回收支持，并且还需要修改链接器以链接模块数据。

作为一个小型团队，我们无法承担如此昂贵的解决方案。因此，我们尝试了不同的方法。

章节 3： 核心方法概述

上一节我们分析了直接实现完整后端的挑战，本节中我们来看看TangoLLVM的核心方法。

我们的方法是将LLVM简单地用作一个后端。以下是工作流程概述：

  1. 首先，运行Go编译器生成Go对象文件（.o文件）。
  2. 接着，在通用的SSA（静态单赋值）级别将Go程序翻译成LLVM IR。通过这种方式，我们可以重用Go编译器中的一些优化，例如逃逸分析。
  3. 然后，运行LLVM后端来优化函数、生成代码，并以ELF格式接收结果。
  4. 我们从ELF文件中提取符号数据，并将其修补回Go对象文件中。

通过这种方式，我们可以避免在LLVM中支持复杂的Go对象文件格式。由于这仍是一个早期项目，我们可能无法使用LLVM成功构建某些函数。在这种情况下，我们可以回退到Go编译器来生成代码。

我们认为，垃圾回收是导致我们无法构建某个函数的主要原因。

章节 4： 支持Go垃圾回收

上一节我们介绍了整体流程，本节中我们深入探讨实现中最具挑战性的部分：在LLVM中支持Go垃圾回收。

Go运行时可以在栈和堆上分配对象。这意味着垃圾回收器必须扫描栈和堆来查找存活对象。

Go的垃圾回收机制有些不同。它不移动或复制堆对象，但可以复制栈对象。这意味着我们必须识别栈上的所有指针。

垃圾回收发生在一个称为“状态点”（state point）的地方。状态点可以看作是一条调用指令及其关联的栈映射（stack map）。栈映射描述了栈的布局，它帮助垃圾回收器找到所有指针。

关于Go ABI的一个事实是：在Go ABI中，没有调用者保存的寄存器。这意味着在状态点，所有寄存器中的值都会被溢出到栈上。因此，我们不需要为支持Go垃圾回收而构建寄存器映射。

以下是我们如何在LLVM中实现垃圾回收支持：
我们简单地基于LLVM中现有的gc.statepoint指令。基本上，我们在IR级别识别指针，因为我们拥有丰富的类型信息。我们在代码生成过程中保持这些信息，并最终在LLVM IR中发出栈映射。

最具挑战性的部分是，LLVM中许多现有的优化过程并不感知垃圾回收。例如，循环不变代码外提（LICM）可能会将指针移出循环，但在移出的位置，该指针可能尚未准备好被引用，这可能导致垃圾回收器出错。

因此，我们尝试在LLVM IR级别识别所有的GC指针。

章节 5： 指针识别与处理

上一节我们提到了识别GC指针的挑战，本节中我们详细看看如何识别和处理不同类型的指针。

我们基本上将所有指针类型的右值视为内部指针（derived pointer）。

  • 内部指针：被垃圾回收器用来定位包含它的对象，然后扫描该对象中的所有指针。
  • 值指针：存活对象内部的指针。
  • 无效指针：对象末尾的指针不是有效指针。例如，如果对象O2的末尾是有效的，它实际上是对象O3的开头，我们应该标记O3而不是O2。

那么，如何识别无效的GC指针呢？我们发现只有内部指针可能需要处理无效指针。一个无效的内部指针可能有一个无效的基指针（例如，空指针），或者其偏移量可能超出对象边界，成为越界指针。

为了识别这些指针，我们使用静态分析。我们首先进行保守的静态分析，尝试识别所有指针。如果任何内部指针可能具有非法的基指针，我们将其标记为指针。我们还尝试传播每个指针所指向值的大小。

静态分析非常保守，可能会报告许多无效的GC指针。但幸运的是，我们发现，如果无效的GC指针没有被记录在栈映射中，它仍然是安全的。因此，如果我们发现一些无效的GC指针，我们可能会将其使用点之前移动。这样，我们可以避免在栈映射中记录这类无效指针。然而，这种方法可能会影响某些优化（如LICM）的效果。

未定义指针是支持LLVM中Go垃圾回收的另一个挑战。Go可以将指针转换为整数，然后再转换回来。我们必须跟踪这些值之间的使用定义链。在这个链中，可能包含一些整数算术运算。某些操作可能不会产生有效的指针，例如，两个指针相加的结果就不是一个有效的指针。

我们将这类值标记为“畸形指针”。一旦我们发现这类畸形指针直接出现在栈映射中，我们最终会放弃编译并回退到Go编译器。因为这种情况在实践中很少见，我们不想为此引入额外的复杂性。

章节 6： 性能评估与总结

上一节我们深入探讨了技术细节，本节中我们来看看TangoLLVM的性能评估结果。

我们使用Go1基准测试套件作为评估对象。我们已经支持了LTO（链接时优化）和PGO（配置文件引导优化）。我们这里有四种配置，每种重复运行20次。

最初，结果中的异常数据让我们感到震惊，因为有100%的性能提升。这通常可能是一个错误。但我们仔细检查了编译器的输出，发现LLVM进行了大量的循环优化，例如将普通循环转换为向量化循环并使用向量化指令进行计算。我们认为这个结果是合理的。

在移除异常值后，我们仍然获得了大约28%的性能提升。

一个有趣的发现是，LLVM非常强大，但如果引导不当，它也可能生成效率较低的代码。例如，在time.format基准测试中，我们发现LLVM的循环不变代码外提将过多的值移出了循环，但在那个工作负载中，实际上只有少数值被真正使用。这就是为什么我们必须在LLVM方法中引入指针优化。

总结

在本节课中，我们一起学习了TangoLLVM项目，这是一个为Go语言构建LLVM后端的尝试。我们探讨了其背后的性能优化动机，分析了标准Go工具链的复杂性，并介绍了一种将LLVM作为“补丁式”后端使用的创新方法。我们深入研究了实现中最关键的部分——在LLVM IR级别支持Go的垃圾回收机制，包括识别和处理内部指针、无效指针以及未定义指针。最后，我们看到了初步的性能评估结果，显示该方法有潜力带来显著的性能提升，同时也指出了需要谨慎引导LLVM优化器的重要性。

033：为Clang做贡献

概述

在本教程中，我们将学习如何开始为Clang编译器项目做贡献。Clang是一个庞大且复杂的项目，对于新贡献者来说可能有些令人生畏。我们将探讨多种入门途径，从简单的任务如问题分类和文档编写，到更复杂的代码审查和功能实现。无论您的经验水平如何，都有适合您的参与方式。

入门途径

如果您有兴趣开始接触Clang，有很多事情可以做，具体取决于您的兴趣和舒适度。

以下是几个主要的入门方向：

  • 问题分类：Clang有一个包含大量问题的数据库。即使您认为自己无法解决这些问题，一个非常有帮助的事情就是验证这些问题是否仍然存在。您可以通过帮助分类问题来开始参与。
  • 文档编写：我们的文档并不完美。这是一个很好的机会，让您可以去了解某个标志、某个诊断信息或某个内置功能，并为其编写文档。
  • 标准相关工作：显然，我们在标准方面做了大量工作。但由于您需要阅读标准文档并理解编译器，这个方向对初学者来说可能更具挑战性。

获取帮助的资源

好消息是，我们有很多资源可以帮助新贡献者入门。

以下是可用的资源列表：

  • 办公时间：我每月举办四次，每次一小时的办公时间。欢迎您来提问，例如“我开始了这个工作，但不知道下一步该怎么做”。
  • 在线社区：我们还有Discord、Discourse论坛，您可以在GitHub对话、会议走廊交谈等场合提问。

代码审查的重要性

上一节我们介绍了入门的几种方式，本节中我们来看看代码审查这个核心贡献环节。我认为我们需要更多的代码审查者。

您不需要是专家也能进行代码审查。您不需要是C++专家或标准专家，甚至不需要完全理解您正在审查的特定领域。您可以通过寻找明显的基础性问题来做出贡献。

以下是代码审查中可以关注的一些方面：

  • 测试覆盖：例如，假设有人正在实现新的诊断信息，这些诊断通常有很多选项。您可以简单地检查所有组合是否都被测试到了。通常并没有。您可以要求编写测试来覆盖所有这些诊断选项。这不需要专业知识，但影响很大。
  • 变更摘要：我强烈认为PR上的摘要是最重要的事情之一。摘要应该能让我完全理解您为什么要做这个更改、您做了哪些更改以及我应该看到什么结果。有时专家们会认为某些事情是显而易见的，但并非对所有人都是如此。
  • 思考更多测试：每当您进行代码审查时，都应该始终思考“我可以为此添加什么测试？”。这可以产生很大的影响。大约一半的情况下，当我要求某人添加更多测试时，他们发现了问题或捕获了一个错误。

代码审查的实践建议

代码审查并不超级正式。您不必将自己添加为审查者，也不必担心如果您点击了“接受”，他们是否会合并一个糟糕的更改。您可以只是跳入任何看起来有趣的内容并提供评论，这非常有益，因为我们有越多的眼睛关注这些补丁，我们引入需要日后修复的错误可能性就越小。

快速获得审查非常重要。如果有更多人可以快速进行审查，那就非常重要。因为如果这是您第一次为Clang做贡献，您付出了努力，但在三周内没有收到任何人的反馈，您可能就再也不会回来了。我们正在努力尽快进行审查，但我们需要更多的人来做这件事。

如果您提交了一个补丁，但一周内没有收到任何人的回复，请随时在Discord上提醒我们，或者在GitHub上提及我们，因为我们可能会错过一些事情。请告诉我们：“我的补丁已经一周没人看了，发生了什么？” 我们会尽力查看。

如何找到合适的审查者

回到最初的问题，知道该找谁审查是每个新贡献者都会面临的一个巨大问题。

我们面临的挑战之一是，我们通过GitHub进行的代码审查不会根据更改的文件自动添加审查者。因此，有时知道该找谁是一个挑战。

我们有一个位于Clang仓库根目录的MAINTAINERS文件，它会告诉您对于编译器的这些区域，您应该联系哪些人。这个列表基本上是我们全年都会保持更新的。我们鼓励人们，如果您有经验并愿意站出来成为维护者，那太好了，我们会把您列出来；如果您决定没有时间再维护了，那也没关系，我们会把您移除。您可以随时加入或退出。因此，维护者文件确实是您应该首先寻求帮助的地方。

从代码所有者到维护者

我们最初是从“代码所有者”开始的，这比之前的“狂野西部”状态是一个很大的改进。但我们发现这对个人来说压力太大了，每个代码所有者都不得不做太多事情，感觉责任太重。所以我们转向了“维护者”模式，这是一种不那么正式的模式，每个部分可以有多个维护者。事实上，当我们看到有人在某个领域做了很多代码审查和补丁，表现出积极、有知识、负责任，并且做了很多代码审查，表明他们了解情况时，我们也会把他们拉入维护者行列。因为将知识分散到更多人身上，对社区和我们个人都有好处。

贡献者的心态与责任

我想说一点，我认为这是隐含的但我们没有明确告诉您：您不会搞砸。您可以尝试一些事情，如果某些事情不正确，无论是代码还是审查评论，只要遵循行为准则并且是出于善意，我们不会严厉批评您。我们会与您合作，帮助您解决这些问题。因此，您可以完全放心地跳入一个问题并尝试解决它。如果您弄错了，会有人注意到，我们会纠正它，然后继续前进。

然而，如果您做得好，我们会尝试给您更多的责任。您不必接受，但我们总是感谢您。我们一直在寻找愿意长期为Clang和LLVM工作的人。所以，如果您愿意做这项工作，我们会给您更多工作，我们会用更多工作来回报您。

合并决策与沟通

关于合并决策，这是一个有点模糊的领域。当您准备点击合并按钮时，您必须感觉到那个人确实进行了彻底的审查，并且具备合理的知识水平。您可能希望给其他审查者一个合理的机会来审查。记住Aaron之前的演示中的图表，周六和周日的贡献量是单点线，而其他日子都很宽。请记住这一点。不要在周五晚上提交，然后说“你们已经有48小时了”。请给一个工作日的时间。

我个人在审查我领域之外的内容时，通常会明确告诉对方：“请注意，我给了您批准，但请先让该领域的维护者检查一下”，因为这不是我的领域。当我这样做时，我通常会说给其他审查者48小时，比如到周二，以便让事情更容易些。

另一方面，如果有人一周前批准了您的PR，但没有人合并它，因为您没有合并权限，请告诉我们。我们有一些PR已经获得批准和审查，准备合并，但却在GitHub上闲置，因为没有人注意到应该有人按下合并按钮。所以，当您的补丁准备好时，请告诉我们，我们会为您合并。

我认为更重要的是，每个人都在围绕一个基本点讨论，那就是一切都关乎沟通。所以，如果您对任何事情感到困惑，就请沟通。您真的不会过度沟通。

选择贡献的领域

Clang是一个如此庞大的项目，我个人希望在所有领域都发展一些专业知识，但由于带宽原因，也许我只能发展其中一个。那么，在Clang中，新贡献者应该从哪个领域开始发展专业知识，而目前这个领域没有得到太多关注呢？

这很多时候取决于您对什么感兴趣。如果您真的对常量表达式感到兴奋，那么我们可以找到方法。在一个生产型编译器中，有太多工作要做，我们几乎可以在您感兴趣的任何领域为您找到事情做。

您也可以在Discord上或办公时间问这类问题，比如说：“我对这个东西感兴趣，您能帮我找到一个切入点吗？”

说到常量表达式，新的常量字节码解释器基本上一直靠一位非常、非常棒的家伙Tim在支撑。他可能需要帮助，这是一个会很有帮助的领域。越快转向用于常量表达式的新字节码解释器，特别是现在我们有了反射功能，我们的编译器就会越快，用户体验就会越好。当它发布时，这将是Clang版本中发生的最大的事情。所以，任何能帮助他加速的事情都是好的。他也在Discord上非常活跃，所以您可以问他需要什么帮助，以及您如何提供帮助。

测试工作的重要性

我们也有很多错误报告。其中一些并不难修复。我们经常标记“好的第一个问题”，它们通常很快被认领，但也不总是这样。所以，如果您想深入了解，错误报告也是一个寻找贡献机会的好地方。我们有一个专门的标签“good first issue”，您可以在GitHub上搜索，它会列出我们认为可能是好的第一个问题的事项。正如Shafiq所说，我们通常会尝试解释为什么我们认为这是一个好的第一个问题，比如“您应该去这里更改这个，然后确保编写这个测试”。

在C++方面，我们可能有2000-3000个开放问题。如果您查看它们，有预处理器、词法分析、代码生成等类别的问题，每个类别都有数百个问题，所以您可以选择任何您喜欢的类别。关于测试，大多数问题都没有实现测试，我们只是不知道。所以，只是编写测试并提交，我们可能需要进行一些调试。

您可能已经注意到一个主题：我们需要测试方面的帮助。如果您有兴趣开始，甚至觉得接触编译器不舒服，只是浏览测试目录，注意到“我没有看到很多关于这个属性的测试”，那就太好了，编写那些测试，为此提交一个补丁。这是一个很好的帮助和学习方式。

文档与缺陷报告

文档也总是我们需要的东西。很难找到优秀的技术文档写手。所以，如果您擅长英语，可以提交文档，这很有用。例如，我们有一个名为AttrDocs.td的文件，其中包含我们所有属性的文档。我们有很多属性只被列为“未记录”。所以，您实际上可以只在一个文件中查找“未记录”，然后说“哦，我了解一些关于cleanup属性或gnu模式的知识”，然后编写那些文档，用户会非常感激，维护者也会感激您。

关于验证缺陷报告，这不仅仅是编写代码来修复缺陷报告。在很多情况下，它们已经可以工作了，但我们需要测试。我们需要有人真正坐下来验证并说“这是好的”或者“哦，不，我们有问题”。

功能原型与上游化

看起来许多较大的功能最终都是先在分支中实现，然后才被上游化。我想知道您对此的看法，这是一个好方法吗？

作为一个生产编译器，我们必须非常小心我们发布的内容，因为一旦我们发布了，人们就会使用它，然后我们就必须永远维护它。所以，如果它是一个实验，如果是一个小实验，有时在编译器内部进行是可以的。但如果是一个像C++合约这样广泛的功能，您真的希望探索整个设计空间，而不想因为发布计划而被锁定。所以我认为在分支中做很多工作是合理的。

这样做的缺点是，我们不希望出现一个巨大的代码转储：“这是10万行代码，现在您有了合约”，因为我们无法审查它。所以，当您在分支中做这项工作时，您必须意识到，将其上游化基本上意味着重做它。

但我们发现，从分支到上游的过程有两个巨大的好处。一是最初编写它的人已经在他们的分支中学到了他们将要学的一切。当您尝试在编译器下、在实验性标志下全新实现某个功能时，您还不知道自己不知道什么。所以，让您进来向审查者解释“我正在做这个，我需要这样做是因为所有这些原因”，比不得不回去修复您搞砸的事情要有效得多。正如Aaron所说，您可能会因为我们已经发布了版本而被锁定。

扩展生态系统

Clang已经发展并拥有了更多不同的方面，从Windows支持到CUDA等等。所有这些都有自己的扩展。您对Clang内部和外部的扩展生态系统有什么看法？

扩展实际上是使编程语言有用的东西。这很奇怪，因为您会想，C++本身当然是有用的。但当您深入研究系统头文件等时，是一堆扩展使其工作。因此，扩展对编译器来说确实是一种好处，但它们也是一个主要的成本，因为每次我们添加一个新扩展，您都必须弄清楚它如何与所有标准功能交互，以及它如何与编译器拥有的其他每个扩展交互。所以，我认为我们目前的门槛设置得相当好，但这很困难，因为我们已经这样做了10多年，积累了大量的扩展。这实际上对标准委员会方面来说可能很困难，因为我们的一些扩展会与标准委员会想要做的事情冲突。所以，我们必须对我们的扩展非常小心，并真正确保它们值得实现的成本，也值得占用那个设计空间的成本。

自动化与工具支持

随着向GitHub的迁移，人们正在为代码审查过程等添加各种 slick 的自动化和工具。其他LLVM的语言前端也有一些非常 slick 的集成。作为一个头脑风暴问题，如果您有资源来帮助自动化某些任务，我们可以构建什么样的自动化来改进贡献流程？

在代码审查过程中，有一系列我们作为审查者必须检查的事项清单，其中很多对我来说感觉是可以自动化的，比如“这实际上是否符合编码标准”这类事情。我们在CI流水线中添加了clang-format来帮助人们发现格式问题。我认为我们还可以做其他事情，比如“我注意到您添加了这四个新的诊断信息，但我没有看到任何包含该诊断文本的测试”，或者“您添加了一个新功能，但没有文档”。所以，我认为我们可以添加更多这类东西，来真正帮助我们作为审查者持续遗漏但主要是清单类的事项。

有一件事需要花费大量时间，我不知道我们能否自动化，那就是尝试进行二分查找和尝试进行归约。如果您有时间做这些，实际上您可以贡献的另一件事是进行归约，以找到像5行代码的最小测试用例。如果有办法自动化归约，那将非常有用。

测试套件与质量保证

说到测试，我认为社区中的很多人并不完全了解，我们实际上有多个测试套件。我们有与Clang一起存放在clang测试目录中的测试，但我们也有LLVM测试套件作为一个完全独立的仓库，用于存放许可证与我们不兼容的测试。例如，我们与GPL不兼容，但我们在LLVM测试套件中有GPL代码，因为我们在那里有Gcc torture测试套件。所以，如果您有代码，您觉得“我认为这不合适，运行测试时间太长，可能会导致CI问题”，我们实际上有其他可以放入测试的地方，这样我们仍然可以从中受益，因为我们确实有运行LLVM测试套件的构建机器人，只是它们不会在每个补丁上都运行。

总结

在本节课中，我们一起学习了如何开始为Clang编译器做贡献。我们探讨了多种入门途径，包括问题分类、文档编写、代码审查和测试工作。我们了解了获取帮助的资源，如办公时间和在线社区。我们还讨论了代码审查的重要性、如何找到合适的审查者、贡献者的心态与责任，以及合并决策的考量。此外，我们涉及了选择贡献领域、功能原型开发、扩展生态系统、自动化工具支持以及测试套件的重要性。无论您的经验水平如何，都有适合您的参与方式，Clang社区欢迎并鼓励您的贡献。

034：概述与挑战

在本节课中，我们将学习如何为D矩阵AI加速器构建一个基于MLIR的图编译器。我们将探讨构建此类编译器的动机、所面临的独特挑战，以及我们采用混合编译策略的解决方案。

大家好，我是Sahiium，与我的同事Shaitch一起，我们将分享为D矩阵AI加速器构建图编译器的一些经验。

在开始之前，理解“为什么”至关重要。为什么要做这件事？为什么要构建一个加速器？

有些人可能听说过“内存墙”、“AI壁垒”或“计算饥饿”这些术语。它们本质上是同一回事，意味着传统架构及其最快的HBM内存无法满足现代AI计算的需求。

因此，我们着手构建一种根本不同的东西，它能为我们提供这些计算所需的10倍、20倍更快的内存性能。

这引出了Corst加速器。Corst架构是使用数据流和模块化构建块构建的，这些构建块可以组合成更大的系统。我们使用内存计算技术来实现快速且高能效的矩阵乘法，并通过大量快速、分布式的片上SRAM来增强，这些SRAM用于存储张量参数。

演讲后面需要注意的一点是，每个芯片被分为四个象限，每个象限是一个软件可寻址的、独立的编程单元。

另外需要记住几点。第一，我们拥有相当复杂的内存层次结构，包括多个层级。但它不基于缓存。我们所有的内存对软件都是透明可见的，程序员和编译器的责任是精心编排数据在这些内存上的放置和移动。

第二点，很可能当我们拥有一个容量相当大的模型时，需要将其分布到多个设备上，并采用适当的并行策略来实现。

现在的问题是，我们如何开始为这种具有花哨内存层次结构、奇怪的内存计算架构思考程序表示，同时还要能扩展到多个设备？幸运的是，我们几乎已将所有这些复杂性抽象成一个紧凑的指令集。

Corst ISA包含少量指令，涵盖数据移动、加载存储、向量和矩阵计算、程序控制以及数据重塑。在右侧，您可以看到一个使用这些指令的简单矩阵乘法的示例，我们将在后面讨论。

有了这个背景，当我们开始构建作为更大软件栈一部分的图编译器时，显然会面临一些挑战。对于这个社区来说，构建编译器很困难并非新鲜事。其中一些挑战是我们独有的，一些则普遍适用，总体上它们属于以下三类：

  1. 通过编译器捕获各种算法的难易程度。在机器学习中，这转化为可以编译哪些类型的ML架构。
  2. 我们为开发者或内核作者提供何种控制和能力，以表达他们的优化提示以及任何程序化构造，并能被下游忠实地表示。
  3. 利用所有这些，不仅是编译器技术，还有程序员的意图，我们能在多大程度上利用架构的真正优势并获得其承诺的性能。

我们解决其中一些挑战（以及更多）的方法是采用混合路线。所谓混合，是指我们有一个基于MLIR的编译器栈，并辅以一个包含高性能和用户优化内核的丰富库。

这些内核可以来自像Triton这样的外部方言，我们可以在功能上支持，但Triton的期望与我们的DSL所需的抽象之间存在结构差异，因此我们有其他方法在稍低的抽象级别编写内核。在每种情况下，内核都会被翻译成MLIR的IR或方言，然后输入编译器进行后续处理和优化。

现在，为了深入了解编译器核心部分内部发生了什么，我将交给我的同事。

谢谢Sahiium。我将从我们混合编译管道的概览开始，然后深入探讨前端、中端和后端中一些有趣的组件。

我们从PyTorch模型开始，通过Torch-MLIR将其降低到我们的编译器前端。在前端，我们执行基于SPMD的分区、填充、分块、量化到我们的数值格式以及常量折叠。此阶段的输出是一系列在张量上操作的仿射循环。

然后是中端。在中端，Triton作为一系列仿射循环集成到我们的编译器中。在中端，我们执行生产者-消费者风格的融合、缓冲区化以及向量化，将操作数据块的操作降低为操作64x64虚拟向量寄存器的操作。此阶段的输出是一系列通过DDR内存相互通信的并行循环嵌套。

接着是编译器后端。在这里，我们的自定义DSL降低到DLIR并集成到编译器中。在后端，我们执行SRE和寄存器分配、指令选择、指令调度、内核缝合以及降低到线程模型。此阶段的输出是DMX IO，即我们硬件的汇编代码。

我提到我们做的一个转换是SPMD分区，让我稍微深入探讨一下。我们在StableHLO中执行SPMD分区，从而得到象限级别的对称计算图，之后的编译管道都在这些象限级别的对称计算图上操作。

SPMD分区并非将计算图映射到计算网格的唯一方式，但对我们来说它非常有意义。我们需要支持的最重要模型之一是基于注意力的模型，SPMD分区允许我们自然地将多头注意力模型中不同注意力头之间的独立性和对称性与我们硬件中作为自主执行单元的象限联系起来。

现在让我们通过具体示例来了解一个512x512的矩阵乘法如何通过我们的编译器栈逐步降低。这将有助于说明MLIR提供的丰富抽象如何使我们能够忠实地建模硬件并据此进行优化。

在管道的早期，我们有一个高级的Linalg砖操作，对应单个矩阵乘法操作。这个单一的矩阵乘法操作随后被分解为64个矩阵乘法，每个大小为64x512x64。我们选择按这些维度分块，因为这是我们的矩阵乘法引擎一次能处理的最大矩阵乘法。

这里需要注意的另一件事是，我们也从FP32数值格式降低到了Corst BFP数值格式。

我提到我们的引擎一次可以处理一个64x512x64的矩阵乘法，但在硬件中，这实际上是作为八个不同的64x64矩阵乘法发生的，并且涉及16个不同的向量寄存器，八个用于激活，八个用于权重。在这张幻灯片中，您可以看到我们朝着这种表示形式逐步降低。

我们最终对矩阵乘法操作进行向量化，并使计算中涉及的16个不同向量寄存器在IR中显式化。IR中还显式化了在这八个不同矩阵乘法之间发生的部分乘积归约，这将它们联系在一起，作为单个矩阵乘法的一部分。

最后，我们降低到基于线程的模型，在此过程中，我们将每个仿射并行循环嵌套转换为在一个组内最多八个核心上运行的GPU调用。我们将每个核心建模为线程。这里需要注意的一点是，尽管我们将核心建模为线程，但我们的硬件实际上并不要求我们这样做，我们的核心实际上可以运行不同的工作负载。

在整个降低过程中，我们使用了多种基于方言的优化，这里已经提到。为了节省时间，我将避免详细讨论它们。

无论代码通过我们的编译器走哪条路径，无论是PyTorch、Triton还是自定义DSL，它们最终都会汇聚到DLIR，这是我们自定义的方言，与我们的硬件ISA具有功能对等性。

DLIR模型硬件有序的四元组执行原语。一个DLIR图代表我们硬件上完全调度的工作负载。DLIR是我们退出MLIR的出口，它相当机械地降低到我们的硬件ISA。

总结一下，MLIR使我们能够解决交付ML编译器工具链的一些关键挑战。它让我们能够满足客户当前的需求，其灵活而丰富的方言系统使我们能够忠实地建模硬件并据此进行优化。它还使我们能够构建一个单一的编译器，可以原生支持代码生成和基于内核的方法。

至于我们的下一步，我们将专注于增加操作覆盖范围。

提升性能并构建内核库。就到这里，谢谢大家。祝大家晚上愉快。

本节课中我们一起学习了为D矩阵AI加速器构建基于MLIR的图编译器的动机、核心挑战以及混合编译策略。我们了解了如何利用MLIR的丰富抽象来建模复杂硬件，并通过SPMD分区、渐进式降低等关键技术将高级计算图转化为高效的硬件指令。

035：概述与挑战

在本教程中，我们将学习如何为分布式量子计算构建一个基于LLVM和MLIR的编译器工具链。我们将从量子计算的基本概念和硬件多样性开始，探讨其带来的独特工程挑战，并了解编译器如何作为连接高层编程与底层物理硬件的桥梁。

量子计算：承诺与挑战

量子优势尚未实现，但工程挑战已经显现。本教程旨在介绍我们看到的未来，并展示我们如何利用LLVM和MLIR来支持量子软件工程师和科学家。

这是一个广泛的概述，从量子硬件如何决定需求，到光子学编译器栈如何应对技术挑战，包括时序约束、语言和中间表示的支持。

量子计算机：异构分布式系统

一个核心的工程观点是：量子计算机是一个带有有线边缘设备的异构分布式系统。它利用叠加和纠缠，使用量子比特来处理信息。其承诺是，对于某些问题，提供比经典计算机快得多的解决方案。

然而，边缘设备的实现方式差异巨大，架构范围从固定拓扑到高连接性，再到可重构系统。

硬件多样性：没有“银弹”

回顾历史，量子技术的发展道路并非直接通向硅基技术，而是在硅基技术主导之前，多种技术竞争并塑造了进步。同样，量子计算的成功不仅依赖于更好的量子比特，还依赖于以架构为中心的方法、网络和分布式设计。

一台实用的量子计算机需要数百万个物理量子比特，因此构建它的方式是横向扩展。编译器需要理解目标硬件。

以下是量子硬件方法多样性的简要概览：

  • 超导量子比特：广泛可用的技术，门操作时间快，已演示数百个量子比特。
  • 拓扑量子比特：由微软主导的相关方法，提供更强的抗错能力。
  • 囚禁离子：提供长相干时间和全连接性，但门速度较慢。
  • 中性原子：在室温下工作，通过穿梭原子到新配置来支持灵活的连接性。
  • 硅自旋量子比特：有望实现具有成本效益的制造。
  • 光子量子比特：利用光的量子特性，提供高连接性和易于扩展性。

由此可见，可用的边缘设备不仅不同，而且差异巨大。控制和测量硬件是特定于模态的，这就在具体技术内部塑造了各自的编译器需求。

共同挑战与编译器需求

所有平台的共同点是：精确控制至关重要，底层操作本质上是硬实时的。技术的多样性限制了现有解决方案的复用，我们离量子电路模型越远，满足特定模态要求就越困难。

我们正处于一个量子比特有限、错误率较高的量子时代。虽然量子优势尚未来临，但我们可以预见大规模容错系统将超越经典计算。

但请思考一下数字：商业应用将需要数千个连接的逻辑（即虚拟）量子比特，以及数量级更多的物理量子比特。每种模态都在与自身的扩展和稳定性挑战作斗争，仅仅扩大规模不太可能解决这个问题。这促使我们倾向于分布式模块化设计和横向可扩展系统作为基础要素。

因此，对于我们基于LLVM的编译器来说，这意味着需要专注于分布式计算，并且复用现有解决方案的能力有限。

现有开源项目与局限

该领域已有知名的开源项目。一个量子编译器将电路降级为硬件诊断中间表示、脉冲序列领域特定语言，再到供应商特定的指令。无论我们如何努力将语言、中间表示和目标平台耦合在一起，这些抽象都会泄漏，并且很难将量子模态与技术解决方案（包括MLIR方言）分离开来。

量子电路的建模无助于解决超越网络的扩展细节。而像脉冲序列这样的低级操作在不同模态间存在差异，这使得抽象变得复杂。分布式多设备工作流在MLIR中表达时，模态非常具体，并且有非常特定的约束，再次限制了复用。

总结

在本节中，我们介绍了量子计算的基本承诺和构建实用量子计算机所面临的巨大工程挑战。我们了解到量子硬件具有惊人的多样性，这直接导致了编译器设计的复杂性。核心挑战在于如何为这种异构、分布式、硬实时的系统构建一个高效、可靠的编译器工具链。在接下来的章节中，我们将深入探讨如何利用LLVM和MLIR来应对这些挑战。

分布式量子计算编译器工具链：第2章：架构与约束

上一节我们介绍了量子计算的多样性和挑战，本节中我们来看看一种具体的混合架构及其带来的编译约束。

自旋-光子混合架构

我们的自旋-光子混合架构通过结合用于存储的硅自旋量子比特和用于高连接性的硅光子，来解决连接性挑战，从而为编译器门实现任意的非本地连接。这是一个多步骤的概率过程，具有网络感知能力且资源密集，需要在低级语言中支持原生控制流。

实际约束：模态特定与硬实时

实际约束仍然是模态特定的。操作受相干时间限制，这使得硬实时处理和低延迟至关重要。我们需要闭环控制、嵌入式控制器来满足延迟要求，以及控制器与运行时之间的高效交互，以支持非破坏性的中途测量。

我们模态中实现效率的支柱是可复用模型和高效的软硬件协同设计，以满足纳米级同步控制器的要求。

分布式编译：从全局模型到设备二进制码

有意义的量子计算需要将各个模块联网在一起。因此，从技术上讲，编译器是为分布式嵌入式系统服务的。光子控制系统是一个量子比特控制器网络，编译器必须将全局控制模型转换为具有时序和通信保证的配对设备二进制码。

全连接性是该架构的一大优势，同时由于规模和复杂性，它也是一个巨大的技术挑战。

编译器支持的复杂编程模型

编译器支持一个用于异构计算的复杂编程模型，具有分布式主机-控制器交互，为运行时提供调度上下文（如物理量子比特本身等共享资源）。这种复杂性对编译器工程师来说是个好消息，因为语言和中间表示在编码、表达和暴露这个编程模型方面扮演着关键角色。

总结

在本节中，我们探讨了一种具体的量子硬件架构——自旋-光子混合架构，并分析了它给编译器带来的核心约束：硬实时处理、分布式控制、复杂的资源调度以及网络感知的编程模型。这些约束定义了我们的编译器必须解决的关键问题。接下来，我们将看到LLVM和MLIR如何为构建这样的编译器提供基础。

分布式量子计算编译器工具链：第3章：LLVM与MLIR的基础与鸿沟

上一节我们了解了量子架构的特定约束，本节中我们来看看经典编译器框架LLVM和MLIR如何作为我们工具链的基础，以及它们与量子需求之间存在的鸿沟。

LLVM/MLIR：强大的基础

从头开始创建一个可扩展的量子编译器栈本身就足够具有挑战性。因此，利用像LLVM这样的主流经典编译器框架是一个重大优势。其基础设施、优化和功能为构建工具链提供了强大的基础。

然而，我们需要协调经典编译与量子资源，并针对量子特定指标（如电路深度和门数量）进行优化。因此，高层编程与物理硬件控制之间的鸿沟比在经典系统中要大。

MLIR：支持量子语义与虚拟ISA

在光子学，MLIR支持量子及模态特定的语义，并支持虚拟指令集架构。我们的量子比特控制器使用增加了实时语义的RISC-V。但脉冲级控制对于上游的MLIR方言来说仍然是一个挑战，现有的方言无法捕捉我们的控制参数。

因此，LLVM和MLIR确实提供了许多构建模块，但构建量子软件栈的关键部分在上游是缺失的。

构建量子编译器：初始焦点与挑战

从零开始构建量子编译器，我们最初的焦点是低级控制实验、基准测试和校准库。平衡本地执行效率与模块化组件调度、交互控制的大局观，需要在纳秒精度上进行决策。因此，决策需要在编译时做出。但同时，控制流限制了静态调度，使得实时反馈循环成为硬件协同设计的一个关键挑战。

量子编译器不可或缺的环节，如量子比特映射和错误缓解，在LLVM中并非原生支持，必须作为自定义转换来实现。工具链的这一部分是量子编译器，它与经典编译器紧密交互，以满足由量子比特相干时间决定的时序要求。

总结

本节我们认识到，LLVM和MLIR为量子编译器开发提供了坚实的工程基础，特别是其模块化、多层中间表示的设计哲学与量子系统的分层特性非常契合。然而，由于量子计算的独特性（如硬实时、概率性操作、特定硬件约束），我们需要在它们之上构建大量自定义的组件和优化。下一节，我们将描绘整个编译器栈的完整层次结构。

分布式量子计算编译器工具链：第4章：多层编译器栈

上一节我们讨论了基础与鸿沟，本节中我们来看看为分布式量子计算设计的完整的多层编译器栈是如何组织的。

量子系统的深层分层

量子系统是深度分层的，从高级语言和中间表示、逻辑门和纠错，到物理量子比特、同步原语和控制电子设备。LLVM为此提供了坚实的基础，而MLIR则提供了一个框架，用于集成量子和经典方言以及纠错和模态特定协议，这与量子计算的多层性质相匹配。这对于底层至关重要，在底层，精确的脉冲和通信调度被转换为控制器级代码，并与运行时环境交互。

编译器栈图示

下图展示了一个用于分布式量子计算的多层编译器栈：

  • 前端：将Q#和其他领域特定语言转换为为自旋-光子架构设计的MLIR方言。
  • 运行时集成：定义辅助拓扑，支持原生非确定性量子协议。
  • 中端：适应系统配置和底层物理门定义的变化。
  • 底层：脉冲语言构成量子电路的构建块，并从高级编译器的视角提供一个虚拟ISA。

在这里，单一的中间表示是低效的。多层栈使用可移植的中间表示在逻辑电路、脉冲控制和分布式硬件之间架起桥梁，并在各层之间进行稳定和抽象。这些层是集成的，但又是不同的，具有不同的范围、目标、用户输入语言和约束。

多种用户与统一表示

除了创建量子电路（例如用Q#）的终端用户外，还有两种内部用户：

  1. 量子软件工程师：创建纠错方案，将逻辑电路转换为具体的容错、可执行调度。
  2. 低级语言用户：使用低级语言进行表征、基准测试，并最终开发更好的量子硬件。

我们需要一个编程模型的统一表示，作为量子编译和完全硬件降级之间的锚定点。在量子计算中，有趣的是多层编译器栈为多个可移植中间表示提供了多个接入点（在图中用红色箭头标出）。

关键中间表示

以下是关键的可移植中间表示：

  • QIR：由QIR联盟维护的行业范围成果。
  • 虚拟ISA：编码量子比特控制器网络的执行模型，在栈的各层内表达同步和通信原语。
    • 逻辑虚拟ISA：仍然是抽象的，位于层之间，并非硬性可执行，更像编译器后端的API。
    • 物理虚拟ISA：抽象硬件实现细节，实现不同版本量子控制系统之间的可移植性。

这种方法支持前向兼容性、更快的迭代以及跨层的独立创新。

总结

本节我们勾勒出了一个为分布式量子计算量身定制的多层编译器栈。它通过引入多个稳定的、可移植的中间表示（如QIR和不同层级的虚拟ISA），来解耦高层算法、硬件控制逻辑和具体的物理实现。这种设计提供了灵活性、可维护性和应对硬件快速迭代的能力。接下来，我们将深入其中一个关键环节：脉冲级语言。

分布式量子计算编译器工具链：第5章：脉冲级语言与执行模型

上一节我们介绍了多层编译器栈的整体结构，本节中我们聚焦于底层的关键——脉冲级语言及其执行模型。

模拟与物理操作

区分模拟量子计算和物理操纵量子比特非常重要。模拟是需要的，但精确的量子比特操作才是最终关键。因此，我们优先考虑脉冲语言和ISA。

脉冲级用例与重要性

脉冲级用例包括校准、控制序列设计、保真度基准测试和实验，以推动硬件改进。它类似于逻辑虚拟ISA，是量子电路的构建块。在自旋-光子模态中，脉冲语言内部的原生控制流至关重要，因为其执行模型是概率性的、网络感知的，并且需要控制和测量设备之间的对齐。确保正确的时序、同步和通信对于精确的量子控制至关重要。

编译器实现与关键约束

我们的光子学编译器使用MLIR来编码用户级虚拟ISA和控制器操作，从而实现从逻辑电路到硬件的完整量子执行。编程模型确保全局协调的量子比特操纵。

关键约束包括：

  • 精确脉冲时序：硬件优化和纳秒级时序至关重要。
  • 资源冲突：需要无冲突的资源分配。
  • 中途反馈延迟：不一致的分支延迟可能破坏量子比特对齐，因此我们使用静态调度和确定性分支来保持对齐的相干执行。

同步和通信原语使得在分布式环境中能够进行无冲突的资源分配和一致对齐的操作，但延迟和调度挑战比在经典系统中更为严格。

编译器实现与Python集成

编译器在底层用C++实现，同时也作为Python包发布。Python绑定暴露了完整的翻译、优化、分析API。与Python的集成支持用户管理实验，Python中的数据集成和系统参数（如设备拓扑、校准参数）直接反馈到编译器中。这允许用户和开发者在适当的抽象级别插入编译器，结合了性能与易用性。

总结

本节我们深入探讨了量子编译器栈的基石——脉冲级语言。它直接控制硬件，必须处理硬实时、概率性执行和分布式同步等极端约束。通过静态调度、确定性分支和丰富的同步原语，编译器确保了量子操作的精确性。同时，提供Python接口使得实验和调试对用户更加友好。接下来，我们将讨论如何将高层量子编程语言（如Q#）集成到这个栈中。

分布式量子计算编译器工具链：第6章：集成QIR与未来方向

上一节我们探讨了底层的脉冲控制，本节中我们来看看如何将高层的量子编程语言（以Q#和QIR为例）集成到我们的MLIR-based工具链中，并探讨未来的标准化方向。

集成Q#与QIR的挑战

回到中间表示作为编译器的关键组件，我们支持Q#作为前端之一，使用QIR来表示门、测量、量子比特遥传和结果处理。将QIR与MLIR集成并非易事。MLIR具有结构化控制流和一等公民的语义，而QIR仍然是LLVM IR，需要仔细的映射。微软的软件栈并非基于LLVM，因此集成选项有限：要么将LLVM IR提升到MLIR，要么在降级之前拦截其内部中间表示。依赖未文档化的实现细节具有挑战性。目前，我们将QIR提升到MLIR，编译降级到物理门，再进一步降级到RISC-V以供硬件执行。

更优路径：MLIR作为稳定中间层

一个更具吸引力的方法，是使用MLIR作为Q#和QIR之间的稳定中间层，以提供丰富、灵活的语义。Mojo语言展示了将编程语言与MLIR概念对齐如何使开发者和行业受益。Q#可以遵循类似的道路以获得更广泛的采用。

将QIR提升到结构化控制流再融入MLIR管道，并非匹配Q#、QIR与现代基于LLVM的编译链的唯一问题。QIR使用指针类型和内存语义来建模量子实体，而我们缺少与量子计算自然对齐的SSA值语义（确保值被赋值一次并显式转换）。QIR的中立性是有代价的：难以表达模态特定功能，也难以向编译器传递元数据来指导降级。

推动QIR发展的MLIR框架

我们并非第一个认为MLIR为QIR发展的下一阶段提供了理想框架的人。MLIR在推动QIR采用方面处于独特地位。在QIR之上标准化架构中立的MLIR方言，可能对整个行业来说都是一个有价值的栈，有助于统一原本碎片化的量子计算生态。

平台特定转换与领域特定语言

进入平台特定领域，我们的目标是避免为复杂的条件算法硬编码编译过程。解决方案之一是建立一个用于量子定制化转换和纠错的标准库。这使得我们的用户（量子软件工程师和物理学家）能够用高级语言指定降级逻辑，从而将后端特定关注点与应用程序逻辑分离开来。

Python领域特定语言在这里是一个强有力的工具，因为用状态机编写类似的纠错代码很繁琐。用户熟悉Python，而领域特定语言可以强制执行模态特定的约束。当脉冲序列方言在底层表达逻辑时，这个库使用户能够影响降级，例如，通过描述量子纠错中的“重复直到成功”控制流。

结合领域特定语言与硬件控制

为了结合领域特定语言、高级语言和低级硬件控制，我们必须考虑编译器工具链在自适应量子电路中的执行模型。经典反馈指导后续操作，这就提出了工具链组件如何交互、哪一层驱动执行逻辑的问题。我们的方法是使用与分布式运行时集成的编译语言来处理这种异构环境。因此，类似于HPC，领域特定语言被降级为多个ISA的二进制码以及元数据，生成主机程序和用于量子比特控制器的设备特定内核。

我们的目标是将后端特定关注点与核心应用逻辑分离，为量子软件工程师和实验物理学家提供工具来探索新颖的转换模式。这些抽象必须被降级到MLIR中。

Python动态性与MLIR静态性的调和

尽管Python的动态性、反射和控制流与MLIR基于静态SSA的设计相冲突。一个可行的方法是将Python限制为一个可静态分析的子集，在Python AST上操作，并将其翻译成MLIR，同时单独处理外部调用。该领域已有类似的先例，例如用于混合量子经典工作流的Catalyst装饰器，以及像PySelf这样从受限Python生成MLIR的库。这个设计是一个折中路径，强制执行一个受约束的Python子集。

灵感来源：Mojo与未来愿景

Mojo作为MLIR的前端，以及对混合经典AI工作流的引人注目看法，对我们来说是非常重要的灵感来源，这在概念上与经典-量子工作流相似。我们设想一个未来，经典AI和量子系统在统一的流程中交互并相互受益。

在那个尚未到来的光明未来之前，我们记得，尽管量子系统有些奇特，但它们仍然属于异构计算这个更广泛的范畴。

总结

本节我们讨论了集成高层量子语言（如Q#）的挑战与方案，提出了利用MLIR作为更灵活、更强大的中间层来桥接高层抽象与硬件控制的愿景。同时，我们探讨了使用Python领域特定语言和标准库来赋能用户、实现软硬件协同设计的策略。最后，我们展望了量子与经典计算（包括AI）深度融合的未来。接下来，我们将面对维护这样一个复杂工具链的实际考量。

分布式量子计算编译器工具链：第7章：维护、版本控制与结语

上一节我们展望了技术集成的未来方向，本节中我们来看看维护基于LLVM的量子编译器工具链所涉及的实际挑战，并以此作为本次教程的总结。

维护LLVM基础的成本

最后，LLVM是一个庞大的C++代码库，保持更新需要付出努力，我们需要控制维护负担，但稳定性对我们至关重要。量子硬件实验是科学的前沿，这本身就充满挑战，并且期望软件栈具有稳定性和可预测性。我们降级到RISC-V进行脉冲级控制，这种依赖关系是我们持续考虑的一部分：我们应该自己实现或定制什么，以及我们应该从上游使用什么。如果是使用上游版本，那么版本是什么。

所有权与定制化的权衡

拥有降级管道的某些部分会给我们精细的控制权，这对于应对大规模工程挑战至关重要，但权衡的结果是，我们在自己的代码库中采用变通方法和定制，既不拥有也不定制后端和目标。

关于C++和版本，我们不支持多个LLVM版本以避免复杂性。支持QIR也有这样的代价，但我们宁愿采用变通方法也不使用LLVM 14。因此我们的版本是LLVM 21，并且我们尝试与上游变更保持一致以简化维护。

更广泛的思考

在这最后一张幻灯片上，我想进行更广泛的反思。在17世纪，人们构思了可以代表世界的通用符号语言。几个世纪以来，这些想法仍然是抽象的，但现在有了大语言模型和生成式AI，我们有了先进的实现。量子计算处于类似的位置：前景巨大，不确定性十足，技术困难重重。

经典芯片中的晶体管本质上是完美的。量子比特的两量子比特门保真度在百分之九十几的范围。这个差距定义了我们面前工程挑战的规模。但是，错过实用规模量子计算到来的风险太高，不容忽视。因此，辩论已经转变。问题不再是“是否”会有量子计算机，而是“何时”以及“如何”。

幸运的是，编译器处于这场变革的核心。

总结

在本教程中，我们一起学习了为分布式量子计算构建编译器工具链的完整图景。我们从量子硬件的多样性和独特约束出发，探讨了如何利用LLVM和MLIR的强大基础来构建多层编译器栈。我们深入了解了脉冲级控制、执行模型、高层语言集成（如QIR）以及赋能用户的标准库和领域特定语言。最后，我们也看到了维护这样一个前沿工具链的实际考量。量子计算编译是一个充满挑战但至关重要的领域，它融合了编译器技术、分布式系统、实时编程和量子物理，正等待着编译器工程师们去开拓。

036：利用MLIR编译面向基础的量子编程

在本教程中，我们将学习如何利用MLIR编译器框架，为一个名为“Query”的面向基础（basis-oriented）的量子编程语言构建编译器。我们将了解如何将高级量子算法描述转换为底层的量子电路，并探讨其中的关键编译技术。

概述

当前主流的量子编程语言（如QCL、Q#）通常要求程序员直接思考和构建量子电路，这属于较低层次的抽象。我们开发的“Query”语言及其编译器旨在提供更高层次的抽象，允许程序员以更直观的方式（例如，通过基础变换和经典代码）描述量子算法，并由编译器自动合成优化的量子电路。

语言设计动机 🎯

上一节我们提到了当前量子编程的抽象层次问题，本节中我们来看看一个具体的例子。

在主流量子编程语言中，实现一个特定的基础变换（例如，将四个|+>态替换为带负号的版本）可能需要手动合成复杂的电路。例如，在QCL或现代商业语言中，代码可能如下所示：

// 示例：在传统量子语言中手动构建电路
circuit some_gate(qubit q) {
    H(q);
    // ... 更多门操作
}

而在我们的Query语言中，同样的功能可以简洁地表达为一个基础翻译操作：

# Query语言中的基础翻译
basis_translation |++> -> -|++>

这种表达更接近算法意图，而非硬件细节。编译器负责将此高级描述转换为具体的量子门序列。

编译器架构与MLIR方言 🏗️

上一节我们介绍了Query语言的设计理念，本节中我们来看看其编译器的整体架构。

Query编译器的核心工作是将Query代码（最初解析为Python AST）最终转换为可供现有量子电路优化器和后端使用的量子电路。其软件架构流程如下：

以下是编译器处理流程的关键步骤：

  1. 解析与前端处理：将Query代码（基于Python语法）解析为抽象语法树（AST），然后提取为Query特有的AST（在实现中是一个Rust数据结构）。
  2. 类型检查与宏展开：对AST进行类型检查、宏展开和类型推断。
  3. ** lowering 到 MLIR 方言**：将处理后的AST lowering 到两个自定义的MLIR方言。

我们将重点关注上述流程中的最后一步，即 lowering 到MLIR方言的过程。首先 lowering 到的是 QIRDI 方言。

QIRDI 方言：高层次的量子语义

QIRDI 方言是直接从AST lowering 而来，它镜像了Query语言的语义，可以说是我们MLIR流程中抽象层次最高的方言。这主要体现在其操作是基于量子态的基础（basis）定义的，而非底层的量子门（如NOT、Hadamard门）。

QIRDI方言主要包含两类操作：

以下是QIRDI方言中的两类主要操作：

  • 基础导向的量子操作：例如 state_change（状态改变）、basis_translation（基础翻译）、state_evolution（状态演化）或在某个基础上进行的 measurement（测量）。这些操作都带有指定基础（basis）的属性。
  • 函数式操作：这类操作看起来有点像MLIR中的 func 方言，包含 constant（常量）、call_indirect（间接调用）等。但这里有一个关键的不同点。

那么，为什么不直接使用MLIR内置的 func 方言呢？原因在于Query语言中函数调用的特殊性。

处理特殊的函数调用：反转与谓词化 🔄

在Query以及许多其他量子编程语言中，函数可以以多种方式调用：

  1. 正向调用：正常顺序执行。
  2. 反向调用：以逆序执行函数。
  3. 谓词化调用：函数仅在整个状态空间的某个特定子空间上运行。这是对传统量子语言中“受控”操作的泛化。

更复杂的是，这些操作可以组合。例如，可以获取一个函数的谓词化版本，然后取其逆，再调用它。

在编译器中，我们这样处理这些特殊调用：首先，像 call_rev（反向调用）这样的操作在QIRDI IR中会生成一个函数值。接着，通过一个“具体化”过程，将其转换为带有 reversed=true 属性的直接 call 指令。

真正的挑战在于内联（inlining）这些反向或谓词化的调用。

内联反向调用

在QIRDI方言中，我们假设所有可反向调用的函数（目前）只包含一个基本块。因此，内联一个反向调用就简化为如何反转一个基本块。

我们通过以下步骤实现：

  1. 从基本块的终止符开始，向上遍历操作。
  2. 为每个操作生成其逆操作。这里的新颖之处在于，我们没有为特定操作硬编码此过程，而是定义了一个 ReversibleOpInterface。
  3. 每个实现了该接口的操作都知道如何构建自己的逆版本。我们自底向上收集这些逆操作，然后自顶向下重建出反转后的基本块。

内联谓词化调用

谓词化调用的内联思路类似。我们同样从终止符向上遍历基本块，但这次使用 PredicableOpInterface。对于许多基础导向的操作，谓词化通常意味着修改其基础属性，使其仅在特定的子空间（由谓词指定）上生效。

成功内联后，我们的代码将变成从量子比特分配到测量之间的一条直线型量子操作序列。这对于许多只支持直线型代码的量子硬件来说是理想状态。

量子电路合成 ⚡

上一节我们探讨了如何在QIRDI方言中处理高级语义，本节中我们来看看如何将这些高级操作合成为具体的量子电路。

内联之后，代码虽然已是直线型，但其语义仍是量子层面的（basis-oriented）。下一步是通过方言转换，将QIRDI方言 lowering 到量子电路方言。

我们设计了一个量子电路方言（例如 QuantumCircuit）。它的设计借鉴了前人工作，其特点是量子比特像SSA值一样在不同操作间显式地流动，这使得依赖关系清晰，并且可以相对容易地替换为其他量子硬件厂商的方言。

基础翻译操作的电路合成

让我们以核心的 basis_translation（基础翻译）操作为例，看其合成过程。我们为这类操作合成的电路总体结构包含：一个可逆的经典相位字符串排列核心，外围包裹着处理向量相位（引入或移除负号）的层，以及最外层的标准化/去标准化层（用于移除和重新引入叠加态）。

对于一个具体的两量子比特基础翻译操作，其合成步骤如下：

以下是基础翻译操作的电路合成步骤：

  1. 解包：将输入的2-qubit寄存器解包为两个1-qubit值。
  2. 去叠加：将状态从叠加态中取出（去标准化）。
  3. 施加相位：施加指定的相位（例如π相位，即负号）。
  4. 重新叠加：将状态重新置回叠加态（标准化）。
  5. 重新打包：将两个1-qubit值打包回一个2-qubit寄存器。

整个合成过程在MLIR中通过一个方言转换模式实现，无需复杂的分析上下文。

经典代码的量子电路合成

Query允许编写经典函数（如Oracle），并由编译器合成量子电路。其流程如下：

  1. 经典函数从AST lowering 到一个经典电路方言（如 ClassicalCirc 方言）。
  2. 对该经典电路进行数字逻辑优化。
  3. 使用工具（如借鉴 T 库的方法）将优化后的经典电路合成为量子电路。

性能评估与总结 📊

上一节我们详细介绍了电路合成的过程，本节最后我们来评估一下这种方法的有效性。

我们选取了四个教科书级的量子算法，分别用Query语言和三种常见的面向电路的量子语言实现。针对不同的量子比特数量，我们使用各自的编译器流程生成电路，并送入一个针对未来容错量子计算机设计的资源估算器，来估算其运行时间。

结果显示，对于这些算法，通过Query编译器自动合成的电路，其预估性能与手工编写的电路相比具有竞争力。

总结：本节课中我们一起学习了如何利用MLIR为面向基础的量子编程语言Query构建编译器。我们探讨了：

  • Query语言如何通过基础变换等高级抽象简化量子编程。
  • 编译器如何通过QIRDI和量子电路等多层MLIR方言进行 lowering 和转换。
  • 如何处理反向调用、谓词化调用等高级语言特性。
  • 如何将基础翻译和经典函数等高级操作自动合成为优化的量子电路。
  • 性能评估表明，该编译器能够使高级抽象在性能上与手工编写电路相竞争。

如果您想了解更多细节，可以参考相关论文和开源代码。

问答环节 ❓

问：当处理可逆函数时，对于那些本身不可逆的操作，你们如何处理？

答：我们在AST层面进行类型检查。Query中有两种函数类型：reversible（可逆）和irreversible（不可逆）。如果一个函数被声明为reversible，那么其函数体内所有的操作和子调用都必须是可逆的。目前，可逆函数中不允许包含经典分支等不可逆操作。未来或许可以借鉴可逆经典编程语言（如Janus）的思想来引入受控的可逆分支。

037：使用与改进优化备注

在本教程中，我们将学习LLVM优化备注的当前状态、使用方法以及近期的一些改进。优化备注是编译器向开发者反馈优化决策信息的重要工具，例如哪些优化被应用了，哪些被错过了及其原因。掌握它们有助于我们理解代码性能并指导编译器生成更优的代码。

优化备注简介与动机

优化，特别是向量化等复杂优化，对程序性能有巨大影响。其效果取决于优化是否成功应用。

例如，你编写了一个简单的循环来计算一系列值的范围，然后对结果进行归约。编译器能否成功向量化这个循环，取决于许多因素。这里的关键问题是：我们如何得知编译器是否成功向量化了循环？如果失败了，原因是什么？

一种方法是直接查看生成的汇编代码，但这可能非常繁琐，对于向量化尤其棘手，因为你可能在循环中看到一些SIMD指令，但这并不意味着循环被完全向量化，可能只是发生了SLP向量化。

因此，我们需要更好的工具来完成这项工作。这就是优化备注的用武之地。使用优化备注，编译器可以直接告诉我们某些优化是否被触发。

例如，如果归约计算的是浮点值的总和，优化备注可以告诉我们，需要允许浮点运算重新排序，循环才能被向量化。

另一个例子是，如果归约函数定义在不同的翻译单元中，这可能阻止其被内联。而编译器无法向量化包含函数调用的指令。在这种情况下，优化备注会告诉我们这一点。

优化备注的现状与挑战

优化备注已经存在了一段时间，但由于使用起来仍有些繁琐，我们并未充分利用它们。

首先，大多数优化备注不像我展示的第一个例子那样具有可操作性。其次，你需要知道一些“魔法”标志来传递给编译器，以发现你真正感兴趣的信息。如果你不确定要找什么，就不得不生成所有的优化备注。

这是可以做到的。我们有一个很好的工具叫 opt-viewer，它允许你以HTML页面的形式，将优化备注与源代码交错查看。

问题在于，即使是一个简单的函数，我们也会被各种不同的优化备注淹没，其中大部分对我们来说并非立即可用。

因此，不幸的是，除了小段代码片段，目前使用优化备注来收集实际的性能洞察是相当困难的。

优化备注的设计与分类

优化备注最初被设计为一种与前端无关的诊断机制，用于从优化器和后端向用户提供反馈。

然而，这些诊断更像是遥测数据，而非真正的诊断。因为与错误或警告不同，如果你在代码中看到一个优化备注，通常无法直接“修复”它。它们更多是为了收集编译器正在做什么的数据。因此，与错误和警告相比，优化备注的数量天生就非常庞大。

这也反映在我们有多种不同类型的优化备注上：

  • 已应用备注：表示某个优化被成功应用。
  • 错过备注：表示一个优化机会未被采纳，例如，因为在某些情况下它可能不正确或不安全。这些错过备注是我们最接近可操作诊断（类似于性能问题警告）的备注。
  • 分析备注：这基本上是其他所有类型的备注，例如循环未被向量化的原因、某些启发式决策以及编译器想要告诉我们的各种统计数据。

挑战的另一部分是，优化备注试图满足两种相互冲突的不同用例。

一方面，我们希望将优化备注作为一种教学工具，弥合用户意图与优化器实际能对代码做什么之间的差距。为此，我们需要关于代码中错过的优化机会的低噪音、可操作的备注，以便用户更好地理解哪些优化失败了、为什么失败，并基于这些信息指导编译器获得更好的性能，例如通过传达某些程序不变量、提供特定假设或使用编译指示覆盖编译器可能采用的启发式方法。

另一方面，我们的大多数优化备注是由编译器工程师为编译器工程师添加的，主要用于调试优化过程。在这里，我们确实希望获得所有能拿到的数据。我们希望调整启发式、理解性能回归并评估新优化的影响。特别是，我们希望能够在持续集成中跟踪真实世界项目的优化备注，以便分析编译器在不同版本间的变化，从而标记出潜在的回归。

这里的主要结论是，我们需要改进我们的工具链，以便能够向每个用户展示他们感兴趣的信息。

如何获取优化备注

为了做到这一点，我们首先需要仔细看看如何从编译器中获取优化备注。

优化备注本质上是键值对。它们有一个名称，包含发出该备注的过程名称和函数名称，以及一堆参数（例如字符串）。

我们有三种不同的方式来输出优化备注：

  1. 通过诊断处理器传递回前端，然后前端可以执行任何操作，例如将它们作为诊断信息打印出来。
  2. 将优化备注流式传输到YAML文件以供后续处理。
  3. 一种类似的比特流格式，它也可以让你将优化备注存储在文件中。这种比特流格式比YAML表示更紧凑。

通常每个翻译单元会得到一个优化备注文件，至少Clang是这样做的。所有的序列化基础设施都是LLVM本身的一部分，但需要通过前端标志来设置。在上述每种格式下，你都可以看到如何在Clang中启用它们。

可选地，你也可以使用正则表达式来过滤优化备注，或过滤特定过程的优化备注。这里需要传递的名称是你在过程CPP文件中可以找到的DEBUG_TYPE宏，而不是你从过程管理器使用的过程名称，所以需要小心。

现有工具概览

基于这些不同的格式，我们现在有不同的工具来生成、处理和显示优化备注。我不想详细介绍每一个，但我想让大家了解已经存在的工具，以免我们重复造轮子。

从生成方面看，你可以从Clang、MIR、Swift（可能还有Rust等）生成优化备注。MIR和Swift有点特殊，因为它们也可以为自己的过程输出优化备注，而不仅仅是启用现有的LLVM优化备注。

我认为最成熟的LLVM优化备注，你很可能需要处理的是向量化备注、内联备注和循环展开备注。根据我的经验，这些备注目前有一些非常好的诊断信息。

对于处理、序列化和解析优化备注，这主要基于remarks C++库，它只是LLVM核心库的一部分。但这个库也有一些C绑定。

然后，有一些工具建立在这个库之上。主要的一个叫做llvm-remarkutil，它有多个子命令用于格式转换、聚合等。对于仅限macOS的二进制文件，你也可以使用一个叫做dsymutil的工具。当你合并程序的调试信息时，你也可以用它来将多个翻译单元的优化备注合并到一个文件中。但如前所述，这仅限于macOS。

还有opt-viewer.py，可以用来统计最常见的优化备注；以及opt-diff.py，可以用来比较两个不同版本之间的优化备注文件，它只显示这些文件之间添加和删除了哪些优化备注。重要的是，这些Python工具目前只适用于YAML格式的优化备注。

同样的情况也适用于我之前提到的opt-viewer.py显示工具。当然，你也可以显示由Clang或opt-viewer显示的优化备注。

还有一些其他工具，例如，有一个CI工具叫llvm-opt-report，它是opt-viewer的一个更专业化的版本，只以内联、向量化、循环展开备注的更紧凑形式显示，并且是用C++编写的，所以这个工具也支持比特流优化备注。此外，Compiler Explorer网站上也有一个基于YAML优化备注的工具来查看优化备注。

改进方向：比特流格式与工具

大多数工具都是多年来独立发展起来的，所以目前存在一些碎片化。特别是比特流优化备注没有得到广泛支持，这给使用整个基础设施带来了一些痛点。

因此，作为第一步，我们希望使优化备注成为工具链中的一等公民。我们投入了一些时间来彻底检修比特流优化备注。

主要原因是可扩展性，因为YAML文件会迅速变得异常庞大。对于一个代码库，我们谈论的是GB级别的YAML；如果你从CI中导出多个项目，则可能达到数百GB。而比特流格式的大小仅为YAML格式的10%到20%。

当我们开始时，比特流优化备注有很多问题。主要问题之一是它们只在macOS上得到真正支持，原因是该格式要求将字符串表存储在目标文件本身中。这也意味着我们无法在像opt这样不直接生成目标文件的工具中使用比特流优化备注。

我们对比特流优化备注基础设施进行了重大检修。现在，比特流文件完全独立，不再依赖于目标文件，它们本质上可以作为YAML的即插即用替代品。当你传递优化备注文件时，我们现在会自动检测文件是YAML还是比特流格式，所以现在使用起来容易多了。

比特流文件现在也小了大约40%，并且我们现在可以嵌入更大的二进制块，这在以前是不可能的。例如，你可以用它来在优化备注中存储LLVM位码，以从特定过程中提取重现用例。

为了消除对目标文件的依赖，我们不得不更改启用优化备注的API，因为我们现在需要更仔细地管理优化备注文件的生命周期。因此，在设置优化备注时，你现在会收到一个特殊的文件句柄，你需要保持其存活，直到你发出所有优化备注。你可以在该函数的文档中找到更多详细信息。

新工具：过滤与汇总

现在我们可以实际使用比特流优化备注了，我们需要一些基本工具。例如，我们添加的第一个功能是llvm-remarkutil中的filter命令。顾名思义，它允许你根据函数、过程、优化备注类型和其他条件过滤优化备注。

过滤标志在-help中有文档说明，并且它们都有一个接受正则表达式的版本，用于更复杂的过滤。你也可以使用这个工具在不同文件格式之间自动转换，因为它总是尝试根据文件中的魔数或你传递的文件扩展名（例如，通过-o）来检测正确的格式。

你也可以使用这个工具将多个优化备注文件合并到一个文件中，例如，如果你因为使用Linux而无法使用dsymutil。还有--exclude标志，可以用来反转过滤器，即你不是过滤出你想要的优化备注，而是排除你过滤的优化备注。你还可以使用这个工具来复制和排序优化备注。

为了帮助我们处理来自所有这些不同优化备注的信息过载，我们引入了一个工具来汇总优化备注。这个工具更适用于通用框架，你可以在其中轻松添加新的聚合策略，而无需大量样板代码。

因此，添加新策略就像实现SummaryStrategy接口一样简单，你只需要实现三个函数。你会逐个获得优化备注，并且可以轻松地进行过滤和聚合。目前，我们只实现了一种策略，叫做InlineCalls，用于内联优化备注。

内联优化备注是按调用点发出的，所以它们很快就会变得非常庞大。这个策略通过按调用点汇总内联统计信息来克服这个问题。这里的好处是，拥有这个事后处理工具，允许我们跨翻译单元汇总内联统计信息，从而更容易看到例如在整个程序中，哪些定义不可用，以及内联失败的原因是什么。

这个工具只输出普通的优化备注，所以你可以在运行该工具后，将它们重新输入到opt-viewer中。你还可以获得关于整个程序中成本最低和最高的调用点的信息，以便于调试内联问题。

我们已经使用这个特定功能发现并定位了一个内联问题，其中std::byte被编译成了非常慢的代码。正如我所说，非常欢迎对此做出贡献。实现新的策略非常容易，我认为还有很多机会可以添加新的汇总策略。这是我认为非常容易上手的事情，所以如果有人能向上游贡献这个工具，我将不胜感激。

你还可以选择keep模式，它允许你指定在输出中保留哪些输入优化备注。例如，如果你将--keep=used与inline-calls一起使用，这会保留你的汇总优化备注所基于的所有原始内联优化备注。

高级功能：统计信息与标签系统

检测编译器行为在不同版本之间变化的最佳方法是使用LLVM统计信息，因为统计信息的覆盖范围比优化备注广得多。然而，如果你在统计信息中发现性能回归，很难排查哪些函数是实际原因。

为了两全其美，我们现在可以将每个函数、每个过程的统计信息提取到优化备注中。这是通过透明地替换我们当前使用的STATISTIC宏，使用线程本地实现和一些过程检测技巧来限制优化备注的生成。

我想指出，这个特定功能尚未落地，我们仍然需要开发聚合工具和比较工具，以使其真正强大。但基本的实现已经在一个分支上。例如，你可以看到来自SRA的统计信息，例如有多少加载/存储被移除。

此外，我们正在为优化备注引入一个标签系统。这本质上允许你为每个优化备注分配多个标签，以便根据某些属性对它们进行分组。

例如，我们可以有标签来区分不同的优化备注类型：对编译器用户可操作的，或者对编译器工程师更像统计信息的。你也可以将与某一类优化（如内联、向量化或循环优化）相关的优化备注分组。

整个系统被设计为可扩展的，与优化备注类型不同。因此，优化备注基础设施的用户（如Swift前端）可以引入自己的自定义标签。整个系统与优化备注类型正交存在。之前曾尝试通过子类型化这些其他类型（如分析、错过和已应用）来引入优化备注类别，例如analysis-fp-compute和analysis-aliasing。但这需要大量样板代码，你最终基本上需要添加所有你想要的不同标签的笛卡尔积作为类型。因此，我们决定采用一个更可扩展的、类似标签的系统。

在实现方面，我已经添加了在优化备注中存储标签的支持。

未来展望与总结

那么，接下来是什么？仍然有一些工作需要向上游合并。我当前工作的版本在这个分支上。一些仍然在这个分支上尚未落地的内容包括比特流大小改进、统计信息优化备注以及标签序列化和解析部分。但所有内容都应该在未来几周内落地。

一旦所有内容都就绪，我们就可以将标签暴露给优化备注发射器API，这样我们就可以开始为我们所有的优化备注添加标签。基于标签，我们可以添加类似警告的分组，使优化备注更容易启用和过滤。例如，我们可以有一个像-Rvectorization这样的标签，它只启用所有向量化优化备注，这样你就不必深入源代码来找出需要传递给Clang的DEBUG_TYPE。

希望所有这些工作的最终回报是，我们能够实现一个真正了解标签的出色诊断工具。例如，当某个东西被标记为统计信息时，这个诊断工具可以减去统计信息的值，这应该为我们提供一种非常好的、一次性查看编译器版本间变化的方法。

另一个很好的想法是拥有Python绑定来处理优化备注，这样我们就可以在opt-viewer中获得对比特流优化备注的原生支持。实际上，我在这个分支上已经有了一个可用的绑定原型。所以，如果你检出这个分支并运行opt-viewer.py处理比特流优化备注，它应该可以直接工作。

总结

总而言之，我想说优化备注是弥合编译器与开发者之间反馈循环的重要功能。但我们遇到了一些挑战，尤其是如何提取特定用户需要的特定信息。

在我看来，最好的方法是继续改进基础设施和工具链，并添加更多可操作的优化备注。随着时间的推移，所有这些基础设施都会得到改善。我们在沙箱中拥有所有这些小工具，可以用来分析编译器的性能回归。

到目前为止我们做了什么？比特流格式现在可以作为YAML的即插即用替代品。我们有几个新工具：filter和summary。我们现在可以使用优化备注来按函数、按过程转储统计信息。

展望未来，我希望优化备注标签将帮助我们改进优化备注的组织、过滤和区分基础设施。因此，当我在这里分配任务时，非常欢迎对此做出贡献。这是一个很长时间没有维护的非常大的基础设施。所以，如果我们能聚集一些人，继续为我们的用户（当然也为我们自己）改进开发体验，那将是非常棒的。

038：Arm工具链的采用

概述

在本教程中，我们将跟随Arm编译器工程师Peter的分享，回顾Arm公司如何将其编译器工具链从完全专有逐步迁移到完全基于开源LLVM生态系统的历程。我们将了解这一转变背后的商业、技术和社区驱动因素，以及在此过程中学到的宝贵经验。

章节 1：Arm工具链的起源与专有时代

上一节我们介绍了本教程的主题，本节中我们来看看Arm工具链的起点。

Arm的历史可以追溯到1985年。当时有一家名为Acorn的英国公司。Acorn在1987年发布了一台名为Acorn Archimedes的计算机。Acorn需要为这台计算机开发一个编译器，他们找到了剑桥大学的两位学者Arthur Norman和Alan Mycroft。他们开发的编译器被称为Norcroft编译器。这个工具链用于在Archimedes上开发软件。

随着时间的推移，Acorn的RISC机器部门演变为Advanced RISC Machine，也就是Arm，并从Acorn独立出来。Arm的工具链也必须随之演进。这个早期的工具链，我们暂且称之为armCC。Arm在移动电话和嵌入式领域取得了早期成功。因此，工具链必须被移植到Windows和当时各种Unix系统上以提供支持。

当时的软件开发主要是定制硬件和定制软件，几乎没有开源软件的概念。Norcroft编译器是80年代GCC的早期同代产品。它必须专注于满足当时市场的需求，主要是代码大小优化。如果能节省一点闪存成本，乘以数百万台设备，就能节省大量资金。当时的工具链是商业产品，附带印刷手册和光盘，主要销售对象是公司而非广泛的用户群体。

随着设备变得更加复杂，C++变得越来越重要。Arm当时并不想涉足C++前端业务，因为这是一个相当困难的挑战。Arm当时的优势在于代码生成。因此，他们授权了EDG前端。如果你有幸或不幸在Symbian系统上开发过，Arm的编译器曾是Symbian的平台编译器，而Symbian当时主要使用C++。

这个工具链在当时这个特定市场运行良好。但请注意，当时一切都很分散，没有一个通用的嵌入式平台来推动统一。事实上，直到2004年，Arm才真正有了一个ABI。

章节 2：市场变革与LLVM的机遇

上一节我们介绍了Arm专有工具链的早期发展，本节中我们来看看推动变革的市场力量。

进入21世纪后期，一切开始发生变化。生产定制硬件和定制软件的成本不断上升。在高端市场，人们开始从定制软件转向开放平台，例如Linux。有了这样的平台，你就会得到一个平台编译器，比如GCC。在这种情况下，很难让人们改用其他编译器。因此，一个不能在Linux上编译任何东西的专有编译器在那个特定环境下用处不大，尽管它在嵌入式领域仍然有用。

同时，其他方面也在发生变化。GCC有一个相对较弱的中端优化，但代码生成很强。然而，随着更复杂的处理器出现，代码生成的优势不再像性能那样关键。这时，像GCC这样具有强大中端优化的编译器开始变得重要。

Arm发现自己处于一个境地：在代码生成方面有些落后，市场正在转向开放平台，Arm处理器也变得更强大，开始与其他架构竞争。Arm真正想做的是不再单打独斗。Arm拥有广泛的合作伙伴生态系统，其中许多公司也有编译器工程师。如果能与他们合作，情况会好得多。GNU社区是一个可以合作的领域，但GPL许可证并不适用于所有合作伙伴。因此，还需要其他选择。

大约在2009年，Arm内部首次提到了LLVM。LLVM被视为可以解决armCC许多问题的方案。

以下是当时对LLVM的评估，分为技术和社区两方面：

技术方面：

  • LLVM拥有基于现代C++的代码库，而不是超过25年历史的C代码库，更容易开发和维护。
  • LLVM在中端优化方面很强，而这正是armCC的弱点。
  • 尽管在2009年还处于早期阶段，但LLVM已经提供了一些Arm架构支持，可以进行测试。

社区方面：

  • LLVM的许可证（Apache 2.0 with LLVM Exceptions）更有利，这意味着可以更容易地与其他合作伙伴在代码生成方面进行协作。
  • 其模块化设计让Arm设想：能否将LLVM插入到现有的编译器中？

在弱点方面，技术问题被认为是可以解决的，并非关键。社区方面可能是最大的风险，因为Arm真正想要的是在LLVM中建立一个能够自我维持的Arm社区。虽然事后看来这似乎是必然的，但在2009年，这只是未来可能发生的事情，并非既定事实。幸运的是，最终一切顺利。

章节 3：早期探索与原型设计

上一节我们分析了Arm选择LLVM的原因，本节中我们来看看他们最初的实施构想和早期实践。

基于LLVM的模块化特性，Arm设想了以下几种设计方案：

  1. 替换中端：利用Arm在代码生成方面的优势，同时使用LLVM强大的中端。构想是：用客户源代码输入，转换为LLVM IR，然后转换回armCC的代码生成器。但这涉及到如何编写“桥梁”代码的棘手问题。
  2. 比较LLVM与armCC：在2009年，Clang尚未完成。因此，Arm从一个EDG桥接原型开始。流程是：用armCC开始，然后转换到LLVM，再全程使用LLVM。这提供了一个比较两者的机会。
  3. 利用LLVM的前端：如果armCC的代码生成真的那么好，理论上可以利用LLVM的更多前端（如Clang处理C++代码），然后转换到armCC的代码生成器。

这些构想听起来都很不错，但在实践中却更加困难。

接下来，我们看看在Arm开发原型期间，LLVM社区发生了什么。以下是基于LLVM社区博客的时间线：

  • 2010年1月：Arm开始原型开发。
  • 2010年2月：Clang实现自举（self-hosted）。
  • 2010年：MC（现在依赖的组件）刚刚起步，对Arm的支持还很初级。
  • 2010年：libc++宣布，这对当时没有C++11兼容库的Arm来说非常有趣。
  • 2010-2011年：Arm后端有许多改进，Arm也体验到了从上游合并代码的“乐趣”。
  • 2011年4月：Arm的原型完成，能够编译各种基准测试以与原始编译器比较，并通过了Plum Hall一致性测试套件，这对质量很重要。
  • 2011年：LLVM获得了新的寄存器分配器。
  • 2011年：举行了第一次欧洲用户组会议，可以说是EuroLLVM的前身。

章节 4：经验教训与社区参与

上一节我们回顾了Arm的早期原型和社区发展，本节中我们总结一下从中学到的关键经验。

首先，引用一段2009年LLVM开发者会议圆桌讨论的转述：

“C++ API的易变性是故意的。它允许设计更快地演进，并强烈鼓励所有LLVM用户将他们的改进贡献给项目。任何未贡献的更改都可能在下一个版本中破坏，并增加维护成本。”

Arm在EDG桥接项目中深刻体会到了这一点。他们有两个组件（EDG前端和LLVM），两者都不独立受控，而Arm需要在中间处理阻抗匹配。这显然行不通。

他们还发现，当时LLVM的代码生成能力明显落后于armCC，但这并不意外，且当时无法立即解决。但同时，他们也看到了社区发展的速度有多快，并且在中端优化占主导的基准测试上，两者的差距要小得多。他们看到了未来的趋势。

因此，得出的明确结论是：如果将来要迁移到LLVM，就必须尽可能贴近社区，避免不必要的差异化，以降低未来的维护成本。

同时，Arm也意识到LLVM社区正在迅速成长。就像GNU社区对Arm的Linux生态很重要一样，无论Arm对其专有工具链做什么，LLVM社区都将变得重要。因此，Arm决定加入开源社区。此时，参与的人员并非在专有工具链上工作，而是完全致力于上游LLVM开发。

这给了Arm一个良好的开端，可以增加对LLVM的经验，并尝试在社区中建立影响力或了解社区运作方式。Arm当时已经拥有知名且成熟的汇编器和反汇编器，因此他们开发了一个名为“MC Hammer”的工具，用于进行汇编/反汇编的往返测试，以确保后端编码正确。此外，由于Arm总部在英国，他们也开始赞助欧洲开发者会议，帮助建立社区。

章节 5：战略决策：向Clang迁移

上一节我们讲述了Arm初步参与社区并吸取教训，本节中我们来看看他们做出的关键战略转变。

几年后，Arm发现他们需要同时支持LLVM、GNU工具链（GCC）以及旧的专有工具链。但情况开始变得难以为继。Arm意识到无法以所需的质量水平同时支持三个编译器工具链，必须放弃其中一个。在这个时间点，被放弃的显然将是问题最多的专有工具链。

对于嵌入式系统，可以选择GNU路线或LLVM路线。许可证问题使得选择LLVM几乎没有悬念。Arm也认为，LLVM是比GCC更好的嵌入式编译器基础。因此，他们决定尝试用Clang替换工具链中的独立编译器（armCC）。这是在2014年左右。

以下是当时做出的一些关键设计决策：

  1. 迁移策略：说服用户迁移工具链非常困难。Arm的策略是，不需要在所有方面都更好，但如果在某个细分领域有优势，就可以从那里开始。Arm选择了Armv8架构（特别是AArch64）和C++11库支持作为突破口，这些新特性只在新工具链中提供。这意味着，如果用户要迁移到这些新特性，他们无论如何都需要做支持工作。新项目可以使用新工具链，而其他项目可以继续使用旧工具链，直到新工具链变得更好。Arm预计这个过程需要2到5年。

  2. 开发模型：选择“紧跟上游”（mainline）模型，而不是维护一个定期更新的大型下游分支。主要原因包括：

    • 新的架构和指令需要上游到LLVM，贴近上游使这个过程更容易。
    • 可以实现“上游优先”开发：如果想在开源LLVM和专有编译器中都拥有某个功能，只需在上游实现一次，合并后第二天就可以拉取下来，无需重复工作。
    • 可以获得上游新特性的好处。
    • 缺点是必须持续支付任何下游补丁的维护成本。

  3. 最小化差异：尽可能减少Clang中armCC特有的扩展，尝试通过宏和内联汇编等方式进行模拟，以保持下游补丁数量最少。

  4. 基准测试：对于处理器公司，基准测试（如EEMBC、Dhrystone、Coremark）非常重要。当时在Coremark上甚至有一场“小型基准测试战争”。在这些方面取得好成绩很重要，但相关的优化并不总是容易上游。

  5. 共享二进制工具和库：当时LLD还处于非常早期的阶段，LLVM libc也不存在，现有的二进制工具更像是LLVM开发者的工具，而非面向最终用户的产品级工具。因此，Arm决定与专有工具链共享原有的binutils和库。

最终的结构如下图所示：左侧是armCC和授权的Rogue Wave C++库（不支持C++11），右侧是Clang和libc++。两者共享相同的二进制工具（binutils）和C库。这种设计的负面影响是，这些共享组件必须保持不含LLVM技术，导致存在两个独立的技术基础。虽然这在一定程度上避免了上游的干扰，但也意味着如果上游发生变化，必须在下游做相应的修改，代价高昂。例如，对于链接时优化（LTO），Arm可能是少数实际使用libLTO.so的用户之一，而不是直接导入代码生成器。

章节 6：上游化的挑战与嵌入式平台兴起

上一节我们介绍了Arm向Clang迁移的设计决策，本节中我们深入探讨上游化面临的挑战以及新的机遇。

你可能会问，既然binutils和LLD当时不成熟，Arm为什么不帮助完善它们？或者为什么不尝试将那些下游补丁（如EEMBC、Coremark优化）上游化？

这里需要再次引用之前的观点：任何未贡献的更改都可能在下一个版本中破坏，并增加维护成本。但如果你处于一个细分领域，社区并不总是需要你的更改。上游化不是馈赠，而是请求社区永久承担该更改的维护负担。 上游社区必须能够理解这个更改，并且额外的复杂性必须有充分的理由。

如果是一个像当时嵌入式系统这样的细分领域，而大多数上游开发者都在桌面或服务器领域工作，那么他们就需要付出额外的努力来理解一个他们不熟悉的技术领域，甚至可能一开始就不明白为什么它重要。所以这在那时相当困难。

Arm观察到，当时嵌入式工具链的典型模式是：一个基于Clang的编译器，加上自己专有的链接器、专有的C库、各自为政的代码大小优化、以及用于嵌入式开发的C语言扩展。这方面没有标准，每家都以自己的方式行事，通常缺乏统一原则。也没有一个通用的软件平台来推动人们做同样的事情。如果你想上游化，就必须考虑每个人不同的实现方式，其工作量可能比仅仅做一个小的下游修改高出10倍。因此，人们最终做出了务实的选择，保持下游修改。

然而，当Arm说没有软件平台推动统一时，那只是“尚未”发生。将时间快进到2018年左右，Arm开始看到所谓的“嵌入式软件平台”复兴。就像Linux在桌面端推动发展一样，在嵌入式端出现了像Zephyr、FreeRTOS、Trusted Firmware等开源项目。当然，作为开源项目，它们自然需要一个开源编译器，并以此为基础构建所有基础设施。当时，这个编译器是GNU编译器。如果你不支持GNU编译器接口、GNU链接器脚本和binutils，那么尝试编译这些项目就会失败。Arm开始收到越来越多的支持请求，称无法用Arm的编译器构建某个开源项目。这表明，开源在嵌入式软件中正变得越来越重要，Arm需要采取行动。

同时，Arm也发现，由于他们在LLVM专有工具链上进行了投资，LLVM在某些新款微控制器上的代码生成已经优于GCC。Arm面临着如何将这些优势带给客户的需求。他们需要的是LLVM开源工具链，而不是专有版本，因为开源工具链才能适用于所有这些新兴平台。

章节 7：构建完全开源的LLVM嵌入式工具链

上一节我们看到了嵌入式开源平台的兴起创造了新需求，本节中我们来看看Arm如何利用社区成果构建完全开源的解决方案。

在一个幸运的巧合中，社区已经修复了许多问题，使得一个完全基于LLVM的嵌入式工具链成为可能。Arm特别感谢社区在链接器和二进制工具方面所做的工作，这主要不是嵌入式社区完成的，而是由Sony、Google等公司推动的，目的是让这些工具达到生产就绪状态。

关键进展包括：

  • LLD：在此期间获得了对链接器脚本（linker script）的强大支持。
  • ClangBuiltLinux：Linux内核本身就是一个巨大的嵌入式系统，对工具要求极高。能够构建Linux内核，就意味着有能力构建几乎任何嵌入式系统（只需少量改动）。这是实现嵌入式工具链的关键推动因素。

基于此，Arm创建了名为“LLVM Embedded Toolchain for Arm”的项目。这本质上是一套构建脚本，用于检出LLVM、一个名为Picolibc的开源C库以及（当时的）LLVM libc。LLVM libc当时对嵌入式支持还没有明确计划。Arm将它们组合在一起。

主要弱点是缺乏多库（multilib）支持。多库是指链接器驱动程序根据命令行选项（如架构、调用约定）自动选择对应目录下的库。对于覆盖整个Arm架构的工具链，可能会有数百种多库变体。由于没有多库支持，Arm不得不使用配置文件。初始版本有33个配置文件，用于定位正确的库，这对生产环境来说显然有些笨重，但勉强可用。

构建这个工具链的优势在于，它帮助Arm找到了社区中也在从事类似工作的其他伙伴。Arm可以在FOSDEM、Embedded World 2023等会议上宣传这个开源工具链。当人们在Discourse或邮件列表上询问如何构建时，可以指引他们到这个项目。这是一个良好的开端。

Arm还启动了社区电话会议（大约在2022年），用于协调代码审查和RFC讨论。他们发现许多社区成员都在朝着同一个目标努力。回顾之前提到的嵌入式社区分散且缺乏上游支持的观点，现在情况即将改变。Arm找到了其他愿意审查补丁并推动事项的公司。这是一个非常好的开始。Arm注意到许多其他类似的社区也开始了电话会议，并鼓励处于特定细分领域的开发者查看开发者日历，加入相关的社区电话会议。

章节 8：最终过渡与经验总结

上一节我们介绍了Arm构建开源嵌入式工具链的努力，本节中我们来看最终的决策和整个旅程的总结。

Arm一直在致力于这个LLVM嵌入式工具链，但同时仍需维护专有工具链。他们再次陷入了需要支持三个工具链（GNU、LLVM专有版、LLVM开源嵌入式版）的境地，并且资源不足以同时做好所有事情。因此，必须放弃其中一个。

Arm本质上是在押注：随着时间的推移，这些开源平台将变得比以前的专有接口更重要。 尽管新工具链可能会损失一些代码大小和专有工具链中的嵌入式特定功能，但其关键卖点是：你可以用LLVM直接构建现有的开源软件平台。 Arm希望未来能将LLVM工具链的其他功能（如消毒剂）引入嵌入式领域，这将成为平台向前发展的关键特性。

Arm的长期目标之一是采用LLVM libc作为C库，这样整个工具链就可以从一个代码库构建。

最终，工具链完成了从“Arm Compiler 6”（专有工具链，包含Clang编译器、libc++和专有的binutils）到“完全基于LLVM”（除了Picolibc，LLVM libc目前是可选项）的转变。理想情况下，Arm希望将LLVM libc变为默认选项，而Picolibc作为备选。

以下是Arm在约15年贡献中学到的主要经验：

  1. 上游带宽 vs 下游延迟：如果你能找到与你目标一致的开源社区，他们的发展速度会远快于你。下游开发决策快（低延迟），但上游拥有更多的开发者，发展带宽大得多。
  2. 社区建设：如果社区的发展方向与你的目标不完全一致，你需要进行社区建设。寻找在同一细分领域工作的人，尝试组织定期聚会、协调代码审查。寻找有共同目标的相邻社区。例如，LLVM binutils和LLD链接器脚本的工作最初并非由嵌入式社区完成，而是由ClangBuiltLinux等其他社区为了其他目标完成的。找到共同目标，你仍然可以实现所需的功能。
  3. 社区会演变：当Arm在2014年开始时，感觉像是唯一做嵌入式LLVM开发的。到了2025年，已经形成了一个完整的LLVM嵌入式开发生态。社区会随时间变化。即使你现在是孤身一人，未来也可能不是。保持与社区的联系，尝试引导其朝你想要的方向发展。有时只需要一家公司率先站出来展示成果，其他人就会跟进。

总结

在本教程中，我们一起学习了Arm编译器工具链从完全专有到完全基于开源LLVM的完整迁移历程。我们回顾了其历史起源、市场变革带来的挑战、早期对LLVM的探索与原型设计、从中吸取的经验教训、向Clang迁移的战略决策、上游化过程中的困难、嵌入式开源平台兴起带来的新机遇，以及最终如何构建并转向完全开源的LLVM嵌入式工具链。Arm的故事强调了拥抱开源社区、保持与上游一致、以及通过社区合作解决共同挑战的重要性。这段旅程虽然特定于Arm，但其经验教训对于任何考虑从专有技术栈迁移到开源生态系统的组织都具有宝贵的参考价值。

039：原因与方法

概述

在本节课中，我们将要学习一种名为“按需构建”的C++编译器运行时库部署模型。我们将探讨这种模型的重要性、其带来的优势，以及它对上游运行时库开发提出的新要求。

第一部分：为何采用按需构建模型及其实现方式

在传统的C/C++工具链安装中，运行时库（如libc、libc++、libcompiler-rt等）是在工具链构建时预编译好的。这些预编译的库文件（如.so或.a文件）会随工具链一同安装。

按需构建模型的核心思想是：工具链安装包中包含运行时库的源代码，而非预编译的二进制文件。当用户构建其目标代码时，工具链会根据需要，从这些源代码动态构建所需的运行时库。当然，这个过程会包含缓存等机制，以确保构建“Hello World”这样的简单程序不会耗时过长。

接下来，我们将探讨采用此模型的原因。

采用按需构建模型的原因

以下是几个关键原因：

  1. 支持广泛的目标平台：Clang/LLVM工具链支持非常广泛的目标平台，包括不同的架构、指令集变体、ABI（例如因使用不同消毒器或配置libc++的稳定/不稳定ABI而产生差异）。预编译库通常只支持一种架构和一种ABI。若要让一个工具链支持多种目标，就需要预置大量库文件，并包含复杂的逻辑来选择使用哪一个。按需构建只需一份源代码，工具链可根据需要生成任何变体，这极大地简化了流程。

  2. 应对嵌入式平台的复杂性：嵌入式平台以存在大量重要的微架构变体而闻名。按需构建能灵活应对这种多样性。

  3. 支持运行时配置选项：运行时库本身可能有配置选项。例如，libc++可以配置为包含或不包含文件系统支持；printf可以包含或排除浮点支持以减小代码体积。用户可能还需要为速度、代码大小或基于性能剖析数据进行优化等不同目的构建库。按需构建使得这些定制成为可能。

  4. 参与链接时优化：在构建时编译运行时库，使得它们可以参与应用程序的链接时优化过程。Clang的LTO机制需要IR比特码作为输入，这与预编译库不兼容。此外，内联许多小型运行时函数（常在循环中调用）能带来可观的性能提升。

  5. 支持基于性能剖析的优化：按需构建允许使用应用程序特定的性能剖析信息来构建运行时库，从而更好地指导内联和其他优化决策。

  6. 改善开发体验：对于运行时库的开发者而言，按需构建使得测试变更更加便捷。开发者只需修改源代码，然后重新构建目标程序即可测试，无需重新构建和安装工具链中所有的预编译库变体。

上一节我们介绍了按需构建模型的诸多优势，本节中我们来看看在Google生产工具链中是如何实现这一模型的。

当前的实现方式

目前，我们的实现是建立在标准工具链安装之上的，并且特定于Google基于Bazel的构建系统。实现方式如下：

  • 构建系统将LLVM运行时库添加为C++应用程序的隐式依赖项。
  • 使用-nostdlib、-nostartfiles等选项进行C++编译，从而依赖这些隐式包含的库。
  • 结合自定义的构建文件与上游LLVM仓库中的Bazel覆盖文件来构建这些运行时库。
  • 目前，该模型已成功应用于libc++和LLVM libc，并正在探索未来将其用于compiler-rt和消毒器。

当前实现是附加在标准工具链之上的。未来，或许可以探索如何将其更深度地集成到LLVM/Clang工具链本身中。

第二部分：按需构建对上游运行时库开发的影响

既然我们已经了解了按需构建模型，现在来看看这种使用模式如何影响上游运行时库的开发，以及我们对库开发者的一些期望。

这是一种与传统预编译、预安装模式非常不同的部署方式，因此也带来了新的、不同的需求。

对上游库开发的具体期望

以下是按需构建模型下，对运行时库开发的一些关键期望：

  1. 简单的构建规则至关重要：目前，我们必须以临时方式将运行时库的构建规则整合到我们的构建系统中。像CMake、Bazel这样的构建工具功能强大，但其中的“巧妙”设计可能使得将其移植到另一个构建系统变得非常困难。更可取的是简单、直接的构建规则：例如，一个文件列表、一个编译选项列表，以及几条将它们编译并链接成库的规则。这更容易被其他构建系统采纳。

  1. 倾向于更少的构建规则：这一点起初让我们感到意外，但事后看来又很明显。当构建系统需要为每次用户构建都解析庞大的构建文件（以判断运行时库是否需要重建）时，解析开销会变得显著。例如，LLVM libc的某些构建文件为每个库函数设置了独立的构建规则，这比整个库只有少数几条构建规则要耗时得多。

  2. 将测试相关规则分离：如果测试规则或测试特定的依赖规则与库的构建规则放在同一位置，构建系统将不得不解析大量与构建运行时库无关的内容。将测试及其依赖放在独立的目录中会很有帮助。

  3. 限制构建时所需的工具：运行时库在用户机器上构建。如果构建过程需要额外的工具（例如，用tblgen生成头文件），那么每个用户都需要安装该工具。更简单的方式是使用#ifdef等预处理指令，这样只需要clang即可。

  4. 确保库是自包含、可提取的：LLVM仓库有超过5万个C++头文件和源文件。我们不希望工具链安装包包含如此庞大的源代码。像libc、libc++这样的库是自包含的。但对于性能剖析库等更专业的运行时，需要避免依赖整个LLVM库。理想情况是，将所需的小部分功能（例如一部分AST定义）提取成一个独立、自包含的基础库，放在一个可以单独提取且不会意外引入新依赖的目录中。

  5. 关于生成的头文件：这是一个个人观点。对于预构建的库，从模板生成头文件并安装是可以的。但对于按需构建，生成的头文件是临时的构建产物，调试信息、断言失败信息都会指向它，用户需要找到这个构建产物才能理解错误信息。因此，我认为使用#ifdef等条件编译指令比基于模板生成头文件更可取。

当然，我们讨论的前提是工具链中安装的源代码必须与源码仓库中的内容一致。但这并非绝对。我们完全可以在工具链构建时对源代码进行修改后再安装。例如，可以用tblgen生成目标无关的头文件，然后将其作为“已安装的源代码”。同样，也可以处理项目的构建文件，提取出相关部分进行安装。一个具体的例子是，源码仓库的构建文件可能包含“安装源代码”的规则，而这个规则在已安装的源代码中是不需要的。

未来的可能性与总结

目前，我们的实现是叠加在标准工具链安装之上的。一个有趣的设计问题是：我们是否应该让工具链学会自动按需构建运行时库？例如，Rust语言已经提供了从源代码构建运行时的选项。这引发了诸如“是否应该调用CMake？”、“是否应该在Clang中实现？”等有趣的讨论。答案可能很复杂，或许需要为运行时库开发额外的专用工具。

总而言之，按需构建的核心思想是：我们希望能够从运行时库的源码仓库中，提取出一套有限、简单的源代码和构建规则，并将其随工具链一同安装。我们期望上游库的设计能够促进这一目标的实现。

最后，请注意今天下午在本会议室还有一场由Paul Kirth和Daniel Thornburg带来的相关演讲，他们将讨论将libc++纳入LTO时遇到的一些不显而易见的注意事项。此外，昨天和去年也有一些与LLVM libc相关的演讲值得参考。

总结

本节课中我们一起学习了“按需构建”C++编译器运行时库的模型。我们首先探讨了该模型在支持多目标、定制化、参与LTO、改善开发体验等方面的优势，并介绍了Google当前的实现方式。接着，我们深入分析了这种模型对上游运行时库开发提出的新要求，包括需要简单、少量的构建规则，分离测试依赖，限制构建工具，确保库的自包含性等。这些考虑旨在使运行时库更易于集成到按需构建的工作流中，为工具链和库的开发者带来更大的灵活性和便利。

040：错失的虚拟化机会调查

在本节课程中，我们将探讨LLVM中全程序去虚拟化（WPD）技术当前存在的一些“错失的机会”。这是一个正在进行中的研究项目，我们将介绍其现状、基本原理以及一些初步的实验结果。

🧠 全程序去虚拟化（WPD）回顾

上一节我们介绍了去虚拟化的概念，本节中我们来看看LLVM中实现的全程序去虚拟化（WPD）是如何工作的。

当一个类包含虚函数时，编译器会生成一些全局值。对于我们的讨论，需要了解三个关键部分：

  • 类型信息：存在于代码中，我们只需知道它存在，无需深究细节。
  • 类型信息名称：一个全局字符串，每个类都有一个。
  • 虚函数表：包含指向虚函数实现的指针，以及一些用于重建类继承关系的类型元数据注解。

其中，虚函数表中函数指针的偏移量对我们至关重要。

考虑一个简单的单继承层次结构，包含三个类A、B、C，每个类都有自己的foo函数实现。在这三个类的虚函数表中，相同的偏移量处存放着指向各自foo实现的指针。

WPD的工作机制是：对于代码中每一个虚函数调用点，检查其关联的类型信息和V表偏移量这对组合。算法会查看该类型信息所指类之下的整个继承层次。如果在该层次下只找到一个唯一的函数实现，那么这个调用就可以被去虚拟化（直接调用该函数）；否则，就不能。

🔍 当前算法遗漏的第一类机会

基于上述回顾，我们可以发现当前WPD算法可能遗漏的第一类优化机会。

考虑一个继承层次：类A和类B各自实现了foo，类C继承自A但没有重写foo。如果有一个类型为C的对象进行虚函数调用，调用点处的元数据会指向类A（因为C使用的实现来自A）。当WPD查看类A之下的层次时，会发现A和B两个实现，因此判定无法去虚拟化。

然而，从逻辑上讲，由于调用对象的确是C类型，并且C是“叶子类”（没有子类），其下没有其他实现，所以这个调用应该可以被去虚拟化。当前算法无法做到这一点。

这里引出一个问题：我们能否依赖对象的数据类型（此处是C）来进行判断？虚拟化传递分析已经假定对象的数据类型是可靠的。虽然理论上存在通过C风格类型转换导致行为不同的情况，但这为我们提供了一个新思路：如果我们能跟踪类型转换，或许可以实现更多的去虚拟化。

💡 由内联暴露的第二类机会

跟踪类型转换的想法引出了另一类能带来显著优化机会的场景，一个来自MySQL的具体代码示例说明了这一点。

以下是相关的代码模式：

  1. 声明一个Json_array类型的变量privileges。Json_array是一个叶子类，其父类是Json_dom。
  2. 获取该变量的地址，并传递给一个名为at_clone的函数。该函数的参数类型是Json_dom*（即父类指针）。

关键点在于：如果at_clone函数被内联，编译器就能跟踪到privileges的地址，并回溯到其声明处，从而证明这个对象在构造时就是Json_array类型。这样一来，即使通过父类指针调用，也能进行去虚拟化。

当然，这其中存在需要解决的正确性问题，研究仍在进行中。目前的第一步是评估这种优化的潜在收益。

这里存在一个阶段顺序问题：通常，去虚拟化能为内联创造机会；但在这个案例中，内联反而为去虚拟化暴露了机会。这带来了一个挑战。

📊 初步实验与结果

作为首次尝试，我们进行了一个简单的实验：在LTO（链接时优化）的内联阶段之后，额外增加一轮去虚拟化。这样，编译时内联暴露的机会可以被利用，而LTO时的内联可能暴露更多机会，我们新增的去虚拟化阶段可以捕捉它们。

以下是一些工作负载的初步结果：

  • 对于MySQL，我们观察到了大量新增的去虚拟化机会（括号内的数字即对应上文讨论的特定机会类型）。
  • 虽然性能提升的百分比看起来不大，但对于像MySQL这样成熟且高度优化的代码库，即使获得0.5%的性能提升也并非易事，因此这个结果值得关注。

✅ 课程总结

本节课中我们一起学习了LLVM全程序去虚拟化当前存在的两类“错失的机会”：

  1. 叶子类调用：当虚函数调用来自一个没有子类的“叶子”类型对象时，即使其父类层次中有多个实现，该调用也可能安全地去虚拟化，但当前WPD算法会将其遗漏。
  2. 内联辅助的类型精确化：通过内联传递父类指针的函数，编译器可以追溯并确定指针指向对象的精确派生类型，从而开启新的去虚拟化机会。这揭示了优化器阶段（内联与去虚拟化）之间有趣的相互依赖关系。

初步实验表明，尤其是第二类机会，在真实负载如MySQL中能带来可观的新优化点。这项调查仍在进行中，旨在探索提升LLVM去虚拟化能力的有效途径。

041：我们是否充分利用了MLIR中的TableGen？

在本教程中，我们将探讨如何在MLIR中更有效地使用TableGen。TableGen在MLIR中常被简化为避免样板代码的工具，但我们将展示，编写更丰富的TableGen代码可以带来显著优势。我们将通过具体示例，说明如何利用现有的上游TableGen工具和一些常被忽视的软件工程实践来实现这些好处。

🧠 优势一：减少编译器开发者的心智负担

上一节我们概述了本教程的目标，本节中我们来看看第一个具体优势：更丰富的TableGen代码可以减少编译器开发者的心智负担。

想象你正在开发一个机器学习编译器，并需要创建一个满足以下要求的填充操作：

  1. 该操作执行边缘填充。
  2. 该操作仅填充输入张量的最内两个维度。

以下是你可能编写的代码：

def MyPadOp : Op<"my.pad"> {
  let arguments = (ins
    AnyTensor:$input,
    AnyTensor:$pad_value,
    I64ArrayAttr:$padding
  );
  let results = (outs AnyTensor:$output);
}

这段代码可以工作，但存在一些问题。根据操作要求，我们可以推断出以下不变式：

  1. 输入和输出张量的秩必须相同。
  2. 输入、输出和填充值的元素类型必须相同。
  3. 输入和输出的维度大小必须相同（除了最内两个可能被填充的维度）。

然而，当前的操作定义并未对这些不变式提供任何保证，这导致了几个问题：

  1. 操作的语法与其语义之间的差距扩大，意味着无效操作状态的可能性增加。
  2. 降低过程变得更加复杂，因为我们必须针对多种可能的无效状态进行防御性编程。
  3. 手写IR更容易出错，因为操作定义约束不足，无法充分利用MLIR LSP服务器等开发工具的支持。

现在，让我们看看优化后的TableGen代码：

def MyPadOp : Op<"my.pad"> {
  let arguments = (ins
    AnyTensorOf<[AnyType]>:$input,
    AnyTensorOf<[AnyType]>:$pad_value,
    I64ArrayAttr:$padding
  );
  let results = (outs AnyTensorOf<[AnyType]>:$output);

  let invariants = [
    // 输入和输出张量秩相同
    CPred<"$0.getType().cast<ShapedType>().getRank() == "
          "$1.getType().cast<ShapedType>().getRank()">,
    // 输入、输出和填充值元素类型相同
    CPred<"$0.getType().cast<ShapedType>().getElementType() == "
          "$1.getType().cast<ShapedType>().getElementType()">,
    CPred<"$0.getType().cast<ShapedType>().getElementType() == "
          "$2.getType().cast<ShapedType>().getElementType()">,
    // 维度大小相同（最内两维除外）
    // 注：此约束可能需要用C++验证器实现
  ];
}

通过对比操作定义和不变式，我们可以看到：

  • 输入和输出张量的秩将相同。
  • 输入、输出和填充值的元素类型将相同。
  • 输入和输出的维度大小将相同（最内两维除外）。

有了这些不变式的保证，我们获得了几个好处：

  1. 我们获得了“构造即正确”的保证。
  2. 降低过程保持机械性，因为我们需要处理的无效状态减少了。
  3. 手写IR的错误率降低，因为我们现在可以获得MLIR LSP服务器等开发工具的支持。

以下是确保获得此优势的具体要点：

  • 尽可能预先指定操作不变式，以获得“构造即正确”的保证，并缩小操作语法与语义之间的差距。
  • 使降低过程尽可能机械化。
  • 获得与MLIR LSP服务器等开发工具更好的集成。

这里提到的所有内容同样适用于属性和类型。我还想强调一些有用的上游操作、属性和类型约束的快速链接，它们对编写更丰富的TableGen代码很有帮助。

🔍 优势二：使编译器行为更明显且更健壮

上一节我们探讨了如何减少心智负担，本节中我们来看看第二个优势：更丰富的TableGen代码可以使编译器的行为更明显且更健壮。

想象你正在开发一个机器学习编译器，并需要创建一个满足以下要求的重塑操作：

  1. 该操作根据指定的形状重塑输入张量。
  2. 该操作不改变输入张量的数据和元素数量。

以下是你可能编写的代码：

def MyReshapeOp : Op<"my.reshape"> {
  let arguments = (ins
    AnyTensor:$input,
    I64ArrayAttr:$new_shape
  );
  let results = (outs AnyTensor:$output);
}

这段代码可以工作，但同样存在一些问题。让我们看看操作的不变式：

  1. 输入和输出的元素类型必须相同。
  2. 输入和输出的元素数量必须相同。
  3. 新形状属性的值必须与输出形状相同。

同样，操作定义并未对这些不变式提供任何保证，这可能导致无意义的操作或令人困惑的错误级联。让我们重点关注后一种情况。

我们可能有这样一个重塑操作：我们有一个4x128的张量，要重塑成2x4x64的张量，而我们的输出类型可能是2x4x32（本应与新形状属性匹配，但由于没有约束，它并不匹配）。我们的编译器可能会将此重塑操作降低为一个拷贝操作。我们将4x128的张量沿列维度分成两部分，然后将每个切片存储到输出缓冲区的4x64块中。这个4x64块是由新形状属性告诉我们的，但我们可能是根据输出类型分配的输出缓冲区。因此，在某些时候，我们将执行越界的缓冲区访问，并得到一个与提取切片相关的令人困惑的错误级联。

现在，让我们看看优化后的TableGen代码：

def MyReshapeOp : Op<"my.reshape"> {
  let arguments = (ins
    AnyTensorOf<[AnyType]>:$input,
    I64ArrayAttr:$new_shape
  );
  let results = (outs AnyTensorOf<[AnyType]>:$output);

  let invariants = [
    // 输入和输出元素类型相同
    CPred<"$0.getType().cast<ShapedType>().getElementType() == "
          "$1.getType().cast<ShapedType>().getElementType()">,
    // 输入和输出元素数量相同
    CPred<"$0.getType().cast<ShapedType>().getNumElements() == "
          "$1.getType().cast<ShapedType>().getNumElements()">,
    // 新形状属性与输出形状相同
    CPred<"$new_shape == $1.getType().cast<ShapedType>().getShape()">,
  ];
}

通过对比操作定义和不变式，我们可以看到：

  • 输入和输出的元素类型将相同。
  • 输入和输出的元素数量将相同。
  • 新形状属性将与输出形状相同。

有了操作不变式的保证，我们获得了构造即正确的合理操作，并且避免了令人困惑的错误级联。这里需要提到的另一个重要点是，新形状属性实际上是冗余的，它的值可以从输出推导出来。

以下是确保获得此优势的具体要点：

  • 尽可能预先指定操作不变式。
  • 尽可能消除操作中的重复信息，以避免令人困惑的错误级联，并在上下文明显的地方实现早期失败。
  • 避免为同一信息设置多个数据源。

🧑‍💻 优势三：降低新贡献者的入门门槛

上一节我们讨论了如何使编译器行为更健壮，本节中我们来看看第三个也是最后一个优势：更丰富的TableGen代码可以降低新贡献者的入门门槛。

想象你刚加入一家拥有机器学习编译器的公司，担任编译器工程师，你的第一个任务是为编译器前端添加2D卷积操作。

你开始工作，在代码库中找到了这段TableGen代码：

def SomeOp : Op<"some.op"> {
  let arguments = (ins AnyTensor:$input, AnyTensor:$kernel);
  let results = (outs AnyTensor:$output);
}

这段代码让你产生了几个疑问。如果你没有机器学习背景，你可能甚至不知道卷积是什么。你可能会怀疑这是一个真实的任务，还是已经有人替你完成了。如果你有机器学习背景，你可能会想这是否是一个1D卷积，以及你是否可以重用它来实现2D卷积。

但是，如果在你的代码库中，你找到的是这段TableGen代码呢？

def Conv1DOp : Op<"conv.1d"> {
  let summary = "1D convolution operation";
  let description = [{
    Performs a 1D convolution over an input signal using a kernel.
    In this compiler, it is used as a building block for temporal pattern recognition.
  }];
  let arguments = (ins
    TensorOf<[F32]>:$input,  // [batch, length, channels]
    TensorOf<[F32]>:$kernel, // [kernel_size, in_channels, out_channels]
    I32Attr:$stride,
    I32Attr:$padding
  );
  let results = (outs TensorOf<[F32]>:$output); // [batch, new_length, out_channels]

  let invariants = [
    // 输入和内核的通道维度必须匹配
    CPred<"$input.getType().cast<ShapedType>().getDimSize(2) == "
          "$kernel.getType().cast<ShapedType>().getDimSize(1)">,
    // 输出长度由公式计算得出
    // new_length = floor((length + 2*padding - kernel_size) / stride) + 1
  ];
}

操作名称清楚地表明了它是什么以及它的作用。描述给出了卷积的一般解释，以及在编译器上下文中关于该操作的更具体信息。操作不变式已被明确指定。

现在，有了这些信息，你或许能够回答之前的问题。你对卷积操作有了一个普遍的和编译器特定的理解。你确实注意到你有一个真实的任务，因为这是一个1D卷积操作。根据操作不变式，你将能够推断出你不能（也可能不应该）将其重用于2D卷积，因为2D卷积将有其自己的一套操作不变式。

总而言之，这段TableGen代码清楚地表明了编译器支持的功能。有了这些信息，作为新贡献者，你可以更独立地完成任务。

以下是确保获得此优势的具体要点：

  • 尽可能预先指定操作不变式。
  • 编写详细的操作描述，以本地化新贡献者处理你的操作时可能需要的所有信息。
  • 清楚地表明编译器实际支持的功能集。

这两点都应有助于新贡献者更高效、更独立地使用你的代码库并开展工作。

📝 总结

在本教程中，我们一起学习了在MLIR中更充分利用TableGen的三大优势及其实现方法。

首先，更丰富的TableGen代码可以减少编译器开发者的心智负担。通过预先明确指定操作不变式，我们可以获得“构造即正确”的保证，使降低过程更机械化，并更好地集成开发工具。

其次，更丰富的TableGen代码可以使编译器的行为更明显且更健壮。通过消除操作中的重复信息并强制不变式，我们可以避免令人困惑的错误级联，并实现早期失败。

最后，更丰富的TableGen代码可以降低新贡献者的入门门槛。通过编写详细的操作描述和明确指定不变式，我们可以本地化所有必要信息，并清晰地展示编译器支持的功能，从而帮助新贡献者更独立地工作。

希望这些具体的示例和要点能帮助你在MLIR项目中更好地利用TableGen的强大功能。

042：面向LLVM的AI驱动的超级优化

概述

在本节课中，我们将学习什么是超级优化，以及如何利用人工智能（AI）来辅助超级优化器，从而发现传统编译器可能遗漏的代码优化模式。我们将深入探讨一个名为Iago的AI驱动超级优化工具的设计、工作原理、效果，并从中总结出对编译器开发有普遍意义的见解。

背景：什么是超级优化？

这是关于此话题最常见的问题。让我们尝试解释一下。

超级优化器提出一个简单的问题：一段代码能否被改进？“改进”的定义由你决定，你可以选择任何指标，如性能、代码大小，甚至能耗。它非常灵活。

你可能会想，编译器做的不也是改进代码这件事吗？这里存在一个细微的差别。编译器必须为大量用户编写的程序良好工作，并且时间有限。而超级优化器不受此限制，它可以花费相当长的时间来完成工作。其工作方式与编译器不同：编译器通过一系列优化遍次（passes）并匹配固定数量的固定模式；超级优化器则通过搜索工作。

它搜索大量候选程序，并找到一个有效的、符合你指标的改进方案。

搜索如何工作？

我们不会深入细节，但给定一个程序片段，超级优化器将枚举大量替代方案。最终，它有望找到一个好的候选。然而，为了枚举这些候选，并且为了保证超级优化器的正确性，它必须进入一个验证阶段，使用SMT求解器或类似Alive2的工具。这两者都是昂贵的操作，导致计算量迅速爆炸式增长。

例如，超级优化一个非常有用且流行的技术SQLite3，仅枚举大小为1和2的程序就需要24小时，且是在一台拥有1.8个核心的大型机器上。

动机：为什么需要超级优化？

从根本上说，它针对一个模式产生一个替换。那么，当LLVM中已经存在InstCombine（指令合并）以及其他编译器也有类似的人手编写的优化遍次时，为什么还要做这个呢？

我们目前能放入InstCombine或GCC的匹配器（如PDD）中的，通常是匹配一到五条指令并将其替换为更短指令序列的模式。这在一般情况下效果很好。

但是，它不可能适用于所有可能的程序。你可能有一段对你非常重要的特定代码模式，而它完全没有覆盖。此外，如果你想匹配更长的序列，在编译器中是不切实际的。你不会想在InstCombine中匹配成百上千条指令。

毫不奇怪，这两点常常是相关的：特定模式往往更长，而更长的模式更具体。尝试优化它们是有价值的。

具体示例

让我们看一个具体的例子。这是一大段代码。我们需要费心优化它吗？我选择这个特定例子是因为在超级优化器40年的历史中，每个超级优化器项目和论文都包含了一个类似的例子。让我展示为什么这整段代码可以简化为一条popcnt（人口计数）指令。

有无数种编写popcnt的方法。这可能是更快的方法之一。如果你认识这些数字，它们来自《Hacker‘s Delight》，这是一种分块执行popcnt的方法。

因此，超级优化器可以将这个东西简化为一条指令。

另一个例子，我们最近在NVIDIA CUDA基准测试中发现了很多sll（逻辑左移）指令。它可以简化为一条and（与）指令。这是LLVM中缺失的优化，我打算很快实现它。

有趣的是，这里的第二个例子不是一个单一的优化。它是两个不同模式的伪装，InstCombine的工作列表必须多次处理才能实现。

但这两个例子，都有可能通过传统的基于枚举的超级优化器找到，因为输出相当有限。

引入AI

所以，我为此使用了AI。

如果你满足于枚举一两条指令，完全不需要使用AI，它扩展性很好，每次尝试大约需要30秒到一分钟。

如果你想枚举更大的程序，比如三、四、五条甚至更多指令，枚举法就行不通了。枚举三条指令对我来说需要接近12到24小时。

我们在这里想使用AI的原因是：与其遍历这个巨大的搜索空间，我们想让AI生成几个选项，然后我们评估它们。最终结果仍然是经过验证的，你仍然能得到正确的东西。这使我们能够更深入地探索搜索空间，找到更大的可优化模式。

优化现有模式

这里有一个非常重要的例子，说明找到更大的模式会有所帮助。左边的代码有一个xor（异或）和trunc（截断）以及shift right（右移），右边则稍微简单一些。这是一个非常、非常常见的操作，几乎出现在我们编写的所有代码中。它检查一个浮点数是否为NaN（非数字）。右边是超级优化器找到的、稍微更高效的方法，前提是你能高效地加载大常量等。这就是你期望从超级优化器中得到的那种东西。

希望这提供了足够的动机。接下来我将转向Iago的设计。

Iago的设计：什么有效，什么无效

我们将探讨在构建基于AI的编译器工具时，一些可能想使用的技巧和窍门。

基线方法是：直接要求一个大语言模型（LLM）：“给定这段代码，请优化它。”不幸的是，如果你这样做，它开始胡言乱语。你会得到一句话，如果你告诉它别的，你会得到“你说得对”之类的回复。这在过去几年略有改善，但如果你只想使用普通的LLM，效果仍然不太好。

如何解决？

这里有两个问题。首先，我们希望输出格式正确，即正确的语法。当然，它必须是一个正确的优化，语义必须正确。让我们看看如何解决这些问题。

第一件事是，你需要一个好的系统提示。你解释超级优化的想法，描述可用的操作（无论你想生成什么），列出AI不应该做的事情（这效果不太好），并给出示例。最后一部分非常重要：根据我的经验测试，给出大约10到15个示例能显著改善效果。超过这个数量帮助不大，但少于这个数量肯定不行。

这使得它比上一张幻灯片中的情况更好，但你仍然没有得到正确性保证。

我们该怎么做？

我们有一个计划。你将结果发送给一个解析器（例如llvm-as或任何其他你想生成的格式）。如果格式不正确，你会得到错误信息。LLM神奇地在下一次生成正确格式的输出，这效果很好。所以，如果你有任何验证语法的方法，给它解析器的错误信息，它会起作用。

有时它不起作用。我将展示两个LLVM的例子。

我注意到在许多结果中，它不断重定义值。我尽力在提示中解释什么是SSA（静态单赋值），但没用。所以有时经过几次尝试它会成功，但更好的方法是直接添加一个后处理步骤，手动修复这个问题，这立即大大提升了我们的结果。

另一个问题是，我很难让它理解类型系统。例如，select（选择）指令产生一个一位的结果，而不是32位；select的条件是一个一位的条件，而不是32位。但你可以通过在代码中间插入适当的修复来解决这两个问题。

有了这两个主要调整，我们得到了正确的语法，即格式良好的LLVM IR。

但它仍然不正确。我们可以尝试用Alive2验证。如果结果有效，我们就完成了，找到了一个正确的转换。

如果无效。类似于解析器错误信息，Alive2可以产生一个反例，你把它反馈给LLM。让我们看看会发生什么，它不起作用。

基于Transformer技术的LLM不太适合理解反例这种否定性示例。当我尝试时，它不断重复产生相同反例的结果。没有简单的方法解决这个问题。

一个选择是不断重试。这有效，但效果不太好。不过，这是我们启动Iago时真正采用的解决方案，如果你没有更好的办法，这也不错。

如何做得更好？

我们提到，这些转换必须是正确的。我们不能在编译器或任何此类工具中实现不正确的转换。我们已经使用Alive2来验证结果。那么为什么不也在这里尝试使用它呢？

给定一个错误答案，我们想修复它以产生正确答案。引导我产生这个想法的见解是：即使Iago的结果是错的，它的结构在某种程度上是对的，如果你眯着眼睛看的话。所以它有正确的结构，部分工作已经完成。

让我们看看我们能对此做些什么。这里有一个例子：假设你有 x * 42 + 1。也许这里的常数42和1是错的。所以，你从中提取一个“草图”，这是程序综合领域的术语。这意味着你删除它的一些部分，在这里是常数。

然后你使用SMT求解器来填补空白，也许求解器能保证那些位置上的正确常数。这是一个非常有用的技巧，可以绕过反例不起作用的事实。

深入细节

不过多深入细节，SMT求解器在这里求解的逻辑公式是这样的。如果你不想读，简单来说，我要求求解器给我一个常数，这个常数能产生一个有效的改进。

但你会看到这个公式的前两项是两个量词，而求解器在处理交替量词时效果不佳。解决这个问题的一种算法叫做“反例引导的归纳综合”（CEGIS）。它基本上将其拆分为两个独立的循环，yasm汇编器求解器是这个算法早期的一个良好实现示例。

回顾

我们有一个大致如下的循环：从Iago得到一个LLM的结果。如果需要，修复它。之后，验证它是否良好。

你可能会认为这个循环需要大量尝试，但并非如此。对于像GPT-5或Claude 4.5这样的现代LLM，大约52次尝试就能获得大部分好结果。之后，是一个长尾，也许尝试200次能得到一个额外结果。所以52次尝试对此基本足够。

到目前为止，我们保证了Iago产生的结果是正确且可靠的。但完备性呢？我们不希望它错过优化。给定一段代码，如果存在有效的优化，我们不希望错过它。

我们能保证这一点吗？

答案是否定的，没有真正的方法能做到这一点。让我展示这里的差距有多大。

这里的蓝色条形图是超级优化器（Superopt）找到的结果数量。每一列代表它尝试综合的指令数量。蓝色和红色条形图之间的差距是Superopt和Iago之间的差距。所以，如果Iago必须产生所有结果，它需要填补那么大的差距。这非常困难，没有简单的方法能做到。

但另一方面，如果你看到Superopt或其他超级优化器难以触及的长尾结果，Iago在这方面效果很好。所以，对于0到2条指令，它不那么好，但在此之后效果非常好。

我发现的一个非常有效的技巧是：如果你在搜索的第一级枚举中，将枚举产生的、被剪枝的候选结果也提供给它，结果会略有改善。但同样，没有真正的方法解决这里的完备性问题。这就是为什么如果枚举在可处理范围内，基于求解器的枚举仍然胜出。

最好的部分是，所有这些都适合LLM的上下文窗口。对于现代LLM来说是这样的。

所以现在我们有了一个可靠但不完备的工具。希望一切都相当不错。

让我展示一个有趣的例子，说明它并不完美。这真正展示了编译器的漏洞和问题是如何渗透进来的，我希望你们都觉得这很有趣。

到目前为止，我们期望的是正确的语法，并且必须是一个有效的改进。所以，当我看到它产生这个时，你可以想象我的惊讶。需要明确的是，这完全错了。它看起来有点正确，但完全错了。

这里发生了什么？这是它产生的原始结果。别试图理解它，我马上告诉你错在哪里。它在结果中偷偷加入了一个前提条件。当然，如果你允许它产生自己的前提条件，它就是正确的，它就是这么做的。

这之所以能通过，是因为我们用来验证的解析器甚至没有考虑到结果中可能存在前提条件。前提条件是输入中才有的东西，不是结果。它就这样溜过去了。这是我遇到的唯一一个有趣的此类例子。

结果

现在我将快速浏览一下结果。Iago产生的结果比Superopt略少一些，但其中许多结果是新的，这些很重要，因为这些结果超出了所有传统超级优化器的范围。

因此，研究这些新结果，可能会引导你在LLVM中实现新的、有趣的优化。

这些是最近在LLVM测试套件上运行的类似结果。这个套件用于跟踪LLVM提交的性能，包含许多多源程序。即使在LLVM使用它跟踪编译器性能一段时间后，我们仍然在这里发现了相当数量的新优化。

同样地，对于Gzip，我们也发现了一些数量略低的转换，因为它是一个有30年历史的成熟代码库。

值得注意的是，对于这些Gzip结果，它在代码大小优化方面（-Oz）与Superopt的表现相当，-Oz是Clang中针对代码大小的优化级别。

我们可能想做这一切的原因之一是它比Superopt更快。让我们看看这是如何运作的。如果你让Superopt枚举一条指令，它更快，但Iago比尝试枚举两条指令的Superopt更快。如果你试图让Superopt枚举三条指令，它基本上大多数时候都会超时，因为枚举三条指令是非常困难和耗时的。

另一个大多数人会觉得有趣的点是：Iago发现的优化是否适用于所有后端？答案是否定的。我想指出的一个有趣现象是，x86和ARM64在判断一个优化是否有利可图方面，比它们各自与RISC-V的一致性更高。

我给你举个例子说明为什么。这是一个相当简单的例子：右边的xor指令在x86上产生两行代码，在RISC-V上产生一行，因为RISC-V有立即数xor指令。还有其他原因导致这种分歧，但记住这确实会发生是件好事。

这是x86与RISC-V、以及ARM64之间的分歧细分。需要注意的一点是，两个对角线方块指出的地方是架构之间对优化有分歧的情况。颜色较浅意味着分歧较少。所以，x86和ARM之间的分歧比RISC-V（这里指RISC-V）要少。

经验与启示

我将以从实现中得出的一些启示来结束演讲。

你在这里发现的东西常常是多个转换的伪装。就像这个例子，这是两个转换的伪装：1）合并两个select；2）将一个select和一个or合并为一个select。

另一个启示是，如果我们现在想在LLVM中实际实现这些，它们可能与LLVM期望的规范化形式冲突。这会导致终止问题，调试起来有点棘手。

第三点是，从中挑选正确的子集来实现，以使最终结果比你原有的性能更高。这是一个非常棘手的问题。考虑到成千上万的优化，我们确实需要好的工具来使这更容易处理。如何从中找出那20个能让你的代码快10倍或10%的优化来实施？

总结

最后，我用这张图来结束我的演讲，它展示了AI和形式化方法之间的共生关系：AI擅长生成文本，形式化方法则不然，但它能确保事物的正确性。

我的演讲到此结束。我想我们有足够的时间提问。

问： 非常有趣。用这种方法你能生成的最长序列大概是多长？因为我认为如果你看传统的Ahar或GSO，它们大概能做到7条指令，随机超级优化大概能做到15条。用这种方法你能做到的最大是多少？

答： 37条，来自这里的这个东西。是的，37条。好的，明白了，谢谢。

问： 抱歉，还有一个问题。当你进行有利性分析时，你只是计算IR通过后端降低后的指令数量吗？

答： 是的，对于这项工作，我有两种方法，但图中展示的是汇编指令的数量。我也有一个使用llvm-mca来统计机器成本的流程，但你也可以想办法运行并测量性能。这有点像微基准测试的工作。是的，我认为这是合理的。比如你可以通过exegesis运行它，但那会很困难，你会遇到上下文问题。所以，指令数量可能是一个很好的初步代理指标，就像你举的那个只有mov的例子，如果你不受前端限制，它只是寄存器重命名。所以，是的，谢谢。

问： 对于这里的NVIDIA同事，如果我们能得到llvm-mca对NVIDIA的支持，那将非常有帮助。

答： 所以。

问： 你谈到了正确性。你是否分析了结果的一致性？我的意思是，对于相同的输入，你是否得到相同的输出？

答： 好问题。OpenAI的LLM曾经有一个温度参数，他们最近移除了。有了那个参数，你可以控制一致性，但现在你不能了。所以你只能依赖于提供这项技术的人来决定你想要多一致。但你可以不断尝试一个输入，一旦得到一个好结果，就坚持使用那个好结果。

问： 好的，所以Alive2在处理更大的函数或循环时 fundamentally 受限。那么我们是否只是错过了一些它本可以用循环找到的很酷的优化？或者你对此有什么打算？

答： 是的，你说得对。所以到目前为止，这仅限于Alive2能够验证的范围。我知道你或这个房间里的其他人正在努力让Alive2扩展到循环。所以请好好干。

问： 嗯。好的，谢谢精彩的演讲。在第30张幻灯片，Iago找到了更多我综合出的零操作示例。这是怎么回事？

答： 综合一个常数会使用我简要提到的那个CEGIS循环。它是一个循环，所以你不能让它永远运行下去。你必须为那个循环选择一个最大界限。所以，每当你超过那个界限，Superopt就放弃了。

问： 还有问题吗？据我所知，你使用了OpenAI。所以是GPT-4还是GPT-5？在Claude的情况下呢？

答： 嗯，算是吧。有一个Azure API允许你通过与OpenAI相同的API访问其他LLM。你尝试过不同的LLM吗？

问： 是的，是的，结果有什么不同？

答： 好问题。所以，根据我的经验，目前Claude 4.5相当不错。GPT-5稍差一些，但仍然很好。所有这些都比我们一两年前的情况要好得多。我第一次尝试是在2023年初，现在的情况比那时好十倍。

问： 再次，好问题。是的，我尝试过推理模型。它并没有显著改善结果，但花费的时间要多得多。是的，是的，花费的时间长得多。

问： 你有没有考虑过，也许不用像GPT这样的东西，而是自己制作AI或者训练它，使其更适应这类任务？比如与一些AI研究人员合作之类的。

答： 根据我与专家的交流，这在两年前很有意义。那时，微调确实带来了巨大的不同。现在看来，这已经过时了，不再流行。有人告诉我，只要给它足够的例子，它就能做得很好。我不知道这有多正确，但这似乎是AI社区的看法。

问： 还有最后的问题吗？现在，让我们再次感谢演讲者。

答： 谢谢。

043：引言与概述

在本教程中，我们将学习如何优化IREE编译器，以使其在运行TinyLlama模型时的性能接近llama.cpp的水平。这是一段面向新手的优化探索之旅，我们将从基准测试开始，逐步深入到MLIR层面的调试与优化。

🎯 IREE项目简介

IREE是一个基于MLIR、可重定目标的机器学习程序编译器。它能够接收多种前端输入，例如PyTorch、TensorFlow等，并在不同的后端上执行，例如x86 CPU、ARM CPU和GPU。简而言之，IREE提供了一个运行时和编译系统，用于生成可在这些运行时上运行的字节码。这使得它非常适用于在边缘设备和异构计算环境中运行各种机器学习程序。

我们团队选择IREE作为学习和理解MLIR及LLVM的切入点，主要基于以下几点原因：IREE本质上是MLIR原生的，它是一个端到端的编译器栈。我们认为其规模非常适合新手，并且它包含了大量针对硬件的抽象。总的来说，它是一个能让你全面接触MLIR编译器栈的端到端解决方案。

🚀 基准测试起点

我们首先在TinyLlama模型上对基于MLIR的IREE项目进行基准测试。最初我们尝试了Llama 3.1 8B Instruct模型，但为了更快的迭代速度，后来决定改用TinyLlama。因为我们的主要目标是针对边缘设备进行优化，而我们的研究小组更侧重于机器人和边缘设备领域。因此，如果你也打算在边缘设备上对大型模型进行基准测试，建议使用TinyLlama而非Llama 3.1 8B Instruct，这样可以获得更快的迭代速度。

我们得到的基础基准测试结果显示，IREE处理每个令牌需要23秒，而llama.cpp仅需84毫秒，性能差距超过200倍。这个巨大的差异开启了我们后续的优化探索之旅，其中涉及了大量的MLIR实践和优化工作。

IREE优化教程：第2章：调试与性能分析挑战

上一节我们介绍了基准测试的起点和巨大的性能差距。本节中，我们来看看在优化初期遇到的调试挑战以及如何应对。

在深入优化之前，我们意识到需要对IREE内部的MLIR栈进行大量调试。我们发现，CPU和内核性能分析工具并没有起到太大帮助。这主要是因为IREE项目本身使用了Python绑定，而C++代码栈被隐藏在这些绑定之后。因此，像CPU性能分析和内核性能分析这类工具，由于无法直接透视背后的C++栈，作用有限。

相反，使用诸如perf之类的C++框架分析器和火焰图分析器提供了很大帮助。通过使用perf来定位最慢的代码帧，再结合MLIR的--print-ir-after-all选项（这是一个MLIR优化选项，Google也经常使用），我们得以找到那些导致性能缓慢的MLIR代码块和瓶颈部分。

实际上，对比优化过程中产生的中间MLIR文件，是观察MLIR在优化间变化的唯一方法。这个过程比较困难，因为你需要阅读大量的文本信息。但这就是作为新手深入MLIR调试时所能采取的主要手段。

以下是调试MLIR时的一个重要注意事项：

  • 处理大型常量：机器学习程序中的张量操作通常包含巨大的MLIR常量，这会严重恶化调试体验，因为MLIR文件可能因此膨胀到4GB、8GB甚至更大。
  • 优化技巧：对于MLIR新手的一个建议是，始终使用dense_resource_elements属性来减小IR文件大小。这本质上是将张量常量从MLIR文件中分离出去。IREE项目在其自身的flow方言中也有一个类似的属性，称为named_parameter属性。如果你正在研究MLIR优化或MLIR文件，很可能会找到一个可以用来分离张量的属性或函数。

IREE优化教程：第3章：优化策略与核心发现

上一节我们探讨了MLIR调试的挑战。本节中，我们来看看实际的优化尝试和关键发现。

在优化旅程的最后，我发现我所尝试的MLIR优化方法实际上都没有显著减少推理时间。真正的性能提升来自于针对推理代码本身的优化，而非MLIR编译过程。具体来说，采用分页扩展的KV缓存技术带来了巨大帮助。

这项优化成功地将处理每个令牌的时间从23秒降低到了421毫秒。这是一个非常关键的结果。

下图展示了优化前后的对比，你可以清晰地看到，在优化前，最耗时的函数是一个等待函数。这表明，如果我一开始就进行跟踪分析，可能会更早地发现瓶颈在于内存访问，而非编译优化错误。

  • 优化前：最耗时的函数是一个等待函数，表明存在内存瓶颈。
  • 优化后：性能得到显著提升，处理每个令牌的时间大幅下降。

IREE优化教程：第4章：经验总结与要点

在本节中，我们将总结从这次优化之旅中获得的经验教训和关键要点。

以下是从这次旅程中得出的一些重要收获：

  • 先观察，后优化：如果我遵循了这个原则，本可以节省大量时间。在深入优化之前，充分进行性能剖析和观察至关重要。
  • 编码与预填充阶段的差异：在使用机器学习程序和编译器时，编码（decode）阶段和预填充（prefill）阶段的工作负载特性是不同的。预填充阶段主要涉及GEMM操作，而编码阶段则更多涉及GEMV操作。如果你正在使用任何ML编译器，请务必牢记这一点。
    • 预填充：C = A * B （GEMM，矩阵-矩阵乘法）
    • 编码：y = A * x （GEMV，矩阵-向量乘法）
  • 始终保持IR精简：为了确保调试工作能够顺利进行，请始终想办法保持中间表示的体积尽可能小。

总结

在本教程中，我们一起学习了优化IREE编译器以匹配llama.cpp性能的完整过程。我们从IREE和基准测试的介绍开始，经历了MLIR调试的挑战，发现了通过优化推理代码（特别是KV缓存）而非MLIR编译本身才是性能提升的关键，并最终总结了“先观察后优化”、注意编解码阶段计算差异以及保持IR精简等重要实践经验。希望这篇教程能帮助那些刚刚接触MLIR和LLVM相关项目的新手们。

044：将泛型代码优化到LLVM-IR

在本节课中，我们将学习如何通过识别和合并等价的泛型函数实例，来优化泛型代码到LLVM IR的编译过程，从而显著减少代码体积并提升编译速度。

背景与问题

C++的模板和Carbon的泛型都允许在编译时对代码进行参数化。这样做的主要目的是提升性能和安全性，例如在编译时进行类型检查。然而，这也会导致代码体积增大和编译时间延长。本次讨论的核心就是：我们能否解决这个问题？

一个简单的例子是，当泛型函数的参数类型不同（如 int 和 double），但在降低到LLVM IR后，它们的参数类型可能相同（例如都变成了指针类型 ptr）。在这种情况下，我们是否真的需要两个函数来完成相同的操作？

因此，我们需要解决的问题是：当存在多个函数实例（或特化），其参数类型在源语言中不同，但在生成LLVM IR后具有相同的LLVM IR类型时，我们能否对这些函数进行去重或合并？

为什么这很重要？

优化泛型函数合并主要影响两个指标：代码大小和编译时间。减少冗余函数可以显著缩小最终二进制文件的大小，并加速编译流程中的优化和代码生成阶段。

现有实现概述

当前的实现在Carbon编译器中完成。Carbon编译器的工作流程包括：语法解析、生成前端IR的检查阶段，以及将前端IR转换为LLVM IR的降低阶段。本次讨论的重点就是这个降低阶段。

拥有前端IR是一个极佳的设计，因为它允许我们以单一函数体的形式表示泛型函数，并通过特定的“符号”进行参数化。这样，我们就可以在前端IR中查询特定指令或类型在给定上下文中的更多信息。

实现挑战与考虑

理论上，比较两个函数是否等价似乎很简单：只需逐条比较指令，并查询前端IR来确认它们在当前上下文中是否相同。然而，实际情况要复杂得多。

以下是实现过程中需要考虑的几个关键点：

  1. 函数调用：如果指令中包含函数调用，我需要知道被调用的函数是否也等价。这意味着我必须考虑所有可能的调用链，即需要完整的调用图信息。
  2. 递归：算法需要能够识别递归调用，并判断是否处于递归中，或者是否有先验信息可用。
  3. 接口与实现：即使LLVM类型相同，通过接口调用的不同实现也可能对应不同的函数体，这必须被考虑在内。
  4. 前端特定信息：某些前端特有的概念需要处理。例如在Carbon中，一个接口可以包含“关联常量”，泛型函数中可能声明该类型的常量变量。因此，算法还需要考虑诸如符号表位置等信息。
  5. 性能：为了通过此优化节省总体编译时间，算法本身必须高效，尽可能减少增加的编译开销。

综合以上考虑，确保正确性的核心在于：必须进行完整的调用图分析，考虑所有LLVM类型信息，并涵盖特定前端的所有相关细节。

算法设计

为了平衡正确性与性能，我们选择为每个函数计算一个“函数指纹”。具体做法是，聚合函数的所有相关信息（如函数类型、依赖于特化参数的指令/类型，以及关键的调用图信息），然后使用高效的哈希算法（如LLVM采用的Blake3）生成一个哈希值作为指纹。

算法的高级定义分为两步：

  1. 生成与收集：生成所有函数定义，并收集相关数据，计算并存储其函数指纹。
  2. 合并与消除：执行等价性判断逻辑，消除所有重复的函数。

在第一步中，我们收集的数据包括函数类型、任何依赖于特化参数的指令或类型，并特别仔细地分析函数调用，以获取调用图信息。

第二步的合并逻辑很简单：对于任意两个特化，检查它们的指纹是否相同。如果相同，则它们等价。算法会返回整个调用图中所有等价的函数对，然后我们可以一次性完成所有替换。

指纹等价的检查，除了比较第一步创建的指纹，还会进行两次遍历以解决特定上下文中的等价性问题。

性能测试结果

我们通过自动生成的Carbon代码进行了压力测试，旨在模拟复杂调用图的极端情况。以下是初步的性能结果：

代码大小方面：

  • 对于平衡类型（包括原始类型和指针类型），优化后LLVM IR的大小减少了45%到50%，汇编代码大小和生成的函数数量也相应减少。
  • 如果大部分是指针类型，函数数量减少可达98%。
  • 作为对比，LLVM内置的mergefunc优化pass在此类测试中，仅对约12%的测试用例有影响，平均减少函数数约1%，最大减少约4%。这表明前端分析能带来的优化潜力要大得多。

编译时间方面：

  • 对于平衡类型，执行此优化逻辑本身导致编译时间增加约0.9%，但使得Carbon的降低阶段（生成LLVM IR）时间减少了2%。
  • 更重要的是，在Carbon的优化和代码生成阶段，时间减少了25%到51%。
  • 对于指针类型居多的情况，由于代码量大幅减少，在降低阶段耗时更少，优化代码所需的时间也显著降低。

总结与资源

本节课我们一起学习了如何通过前端IR分析，为泛型函数生成指纹并进行等价合并，从而有效优化LLVM IR的生成。这种方法能显著减少冗余代码，提升编译效率。

该算法已开源。

  • 算法代码：可在提供的第一个链接中找到。
  • 相关文档：详见第二个链接。
    如果你有更多问题或想深入了解，欢迎联系我。

045：规范化中的位置——我们不需要验证器吗？生成有效的DXIL

概述

在本节课程中，我们将探讨如何通过将验证规则直接编码为转换过程，来生成有效的DirectX中间语言。我们将首先了解当前DXIL生成流程的问题，然后介绍一种新的合法化方法，旨在使中间表示本身在生成时就符合DXIL规范，从而减少对独立验证器的依赖。

术语定义

首先，我们来定义几个核心术语：

  • DXIL：DirectX中间语言。它本质上是LLVM 3.7的一个变体。
  • DXC：当前的生产编译器，负责输出DXIL。
  • DXV：验证器，用于检查DXIL模块的正确性和兼容性。

验证的目的与当前流程

验证的目的是确保我们生成的DXIL是合法的，并且与硬件和驱动程序栈兼容。这对于确保着色器能在不同GPU和API版本上正确运行至关重要。

当前，验证发生在代码生成之后、管道的末端。这意味着验证规则必须影响更早的优化过程，这并不理想。当DXC和验证器版本匹配时，我们不会看到代码生成问题。但在使用旧版本验证器时，这会成为问题。

因为验证发生得很晚，我们只能在最后阶段看到错误。维护兼容性变得被动，规则必须与编译器更改手动同步。这迫使合法化规则存在于前端和优化层中。对于每个版本，我们都有可能支持一个DirectX目标，这可能会影响我们SPIR-V后端的代码生成，这并不理想。

此外，我们有一个要求：新编译器必须能够跨DirectX API版本生成有效的DXIL。验证器是我们对硬件合作伙伴的承诺，确保新编译器不会破坏当前和旧的GPU驱动程序。新编译器确实可能破坏旧的验证器，这可能会延迟发布，使得这个过程脆弱且容易出错。我们不希望在现代化进程中延续这种做法。

验证器的作用与局限

那么，当前的验证器具体做什么呢？它检查结构和语义、资源使用以及着色器模型规则。如果一切合法，它会盖上验证通过的“哈希戳”，驱动程序就知道它是好的。

它不做什么呢？验证器不修复任何问题。如果DXIL是错误的，它只会报错，不会进行修复。因此，我们不想依赖DXV作为安全网，而是希望IR本身在生成时就已经是DXIL合法的。

重新构想管道：三步计划

上一节我们介绍了当前验证流程的局限性，本节中我们来看看我们如何重新构想整个生成管道。这是一个三步计划：

  1. 首先，从一开始就不生成非法的DXIL。我们稍后会详细讨论这一点。
  2. 我们将把验证器规则编码为DirectX后端内部的转换。
  3. 我们将仅把DirectX验证器用作最终检查步骤，而不是作为正确性关卡。

这将使我们能够支持跨验证器版本的规则，以应对从前端和优化器接收到的不断演进的LLVM IR。

因此，流程将从 LLVM IR -> DXIL -> 验证器 -> 失败 转变为 LLVM IR -> 合法化 -> DXIL -> 验证器（仅用于哈希戳）。这使得输出可预测，并消除了交互中的隐藏规则。

数据转换

这些变化中我们必须做的一项就是数据转换。这里我想传达的关键信息是：我们将机械地重建LLVM数据类型和布局，以使其能够被DXIL处理。

我们正在使用现有的标量化路径，将向量操作转换为DXIL支持的标量操作。这在Shader Model 6.8之前是必须的（6.9显然会有向量支持，但在此之前没有）。我们还必须构建自己的数据标量化来处理诸如别名和全局变量之类的东西。你可以看到我们有一个类似这样的转换。

DXIL还有一个要求，即数组必须是扁平的。因此，我们将线性化多维数组。这实际上必须在标量化之后进行，因为当你运行标量化时，向量本质上会变成一个二维数组。

你可以看到我们展示了该转换是如何工作的。

IR转换

接下来我们需要做的是一组IR转换。

我们引入了一个自定义的合法化过程。我们不能为此使用Global ISEL，因为我们有一个要求：我们仍然能够进行位码序列化。没有方法可以从Machine IR回到LLVM IR位码，但这个想法很大程度上受到了Global ISEL合法器的启发。

因此，这个过程会重写任何剩余的不映射到合法DXIL的构造。这实际上是一个“包容性验证器”，它进行转换，而不是拒绝。

那么这具体是怎样的呢？我们有以下几种指令转换：

  • 指令转换：LLVM 3.7不支持 fneg 指令，也没有 freeze。因此，我们必须用合法的等效构造替换它们，或者完全移除它们。
  • 类型限制：DXIL对类型有限制。一个例子是对整数宽度的限制。对于索引以及插入和提取操作，我们完全不支持任何i1类型，它们只能是32位。对于仅作为索引类型的情况，这很简单，我们可以移除或用32位类型替换那些64位类型。当我们试图合法化i1时，情况会变得更复杂，这实际上需要我们遍历使用链并找到类型转换。找到转换后，我们本质上必须选择最小的转换并进行变换，这样我们就不必在之后修改任何更大的转换。你可以在这里看到一个例子，我们有两个冲突的转换，我们选择较小的那个，并将其传播到别名。
  • 处理内部函数：我们通常通过内部函数展开来实现。DXIL不支持任何内存内部函数。这意味着像 memcpy 和 memset 这样的操作需要被展开成符合DXIL内存模型的合法大小形式。
  • 特殊类型处理：我们通过目标扩展类型使用了一些特殊类型。这些类型不能被像 SimplifyCFG 或我们稍后会谈到的 GVN 这样的过程修改。这些资源类型不能流经Phi节点或选择指令。我们需要一种方法来标记它们，以便优化过程不去动它们，并完全阻止某些优化。因此，我们提出了“类令牌”这个概念，它利用了优化过程中已经存在的关于令牌的一些规则。在这里我们可以看到，我们不将其传播到不同的基本块，我们将其保持在不同的基本块内，我们不试图将其下放到一个Phi节点中。

类似地，我们不得不对GVN过程做同样的事情。我们必须教会它什么是“类目标类型”，这样它就不会为其生成文件。这里的想法是，我们只是试图防止其被错误地重构。

总结

本节课中我们一起学习了如何通过将DXIL验证规则直接编码为LLVM IR的转换过程，来从根本上改变有效DXIL的生成方式。这种方法不打算一夜之间取代验证器，但它形成了一条路径，使得每一条验证器规则都可以变成一条转换规则，并确保未来生成的代码以及针对旧版DXIL的代码都能保持兼容。这使生成过程更具可预测性和鲁棒性。

是的，非常感谢。

046：指令成本建模 - 我们能做得更好吗？🚀

在本节课中，我们将探讨 LLVM 中的指令成本模型。我们将了解它的工作原理、当前的优势、面临的挑战，并讨论一些潜在的改进方向。成本模型是 LLVM 优化器做出决策的关键工具，理解其现状和未来对编译器开发者至关重要。

成本模型简介

上一节我们概述了课程内容，本节中我们来看看 LLVM 成本模型的基本概念。LLVM 的指令成本模型是一个简单快速的近似工具。它接收 LLVM IR 作为输入，并为给定的后端提供一个指令成本估算。

它是一个高层次、可调优的启发式模型。它主要被视为一个相对指令成本。最初它试图表示指令的倒数吞吐量，但现在它更多地成为特定后端上指令一般成本的代理。它被用作优化决策的指导工具。

一个常见的问题是：成本模型必须完全精确吗？实际上，成本模型的主要目的是引导优化器做出正确的决策。只要它足够接近，通常就足够了。因此，讨论的焦点往往是：当前的成本模型是否“足够好”。

成本模型的工作原理

成本模型通过目标转换信息（TTI）钩子实现。你向特定后端的 TTI 传递一条指令，它会返回一个成本值。然后，优化器会为正在评估的代码段（例如一个循环或一个基本块）累加这些成本，得到一个总体成本值。

以下是其工作流程的简化描述：

  1. 优化器需要做出决策。
  2. 它生成一个或多个可能优化方案的示例 IR。
  3. 查询 TTI 获取每个方案的指令成本。
  4. 累加成本，并根据总成本做出最终代码生成决策。

当前成本模型的优势

在深入探讨挑战之前，我们先看看当前成本模型做得好的地方。它已经作为一个可行的启发式方法运行了很长时间。

以下是其主要优势：

  • 快速：查询是常数时间复杂度（O(1)），不依赖于代码生成。
  • 可组合：可以累加不同指令的成本，为更高级别的对象（如基本块）计算总成本。
  • 可覆盖：不同的后端可以覆盖默认实现，提供特定于硬件的成本。
  • 共享基础设施：它是 LLVM 中所有后端共享的通用基础设施。
  • 决策工具：主要用作决策指导，不要求 100% 精确。

成本模型面临的挑战

尽管当前模型有其优势，但在我们的工作中也发现了若干挑战。上一节我们介绍了它的优点，本节中我们来看看它面临的问题。

以下是我们在研究中观察到的五个主要挑战：

  1. 可维护性：成本模型基于启发式，代码遍布 LLVM 约 10 万行。不同的优化器（如循环向量化器、SLP 向量化器）以略有不同的方式使用它，实现各自的成本累加逻辑，甚至存在多个成本模型（如遗留模型和 VPlan 模型），这使得同步更改非常困难。

  2. 启发式调整的滞后性：成本模型最初可能代表指令的倒数吞吐量，但已演变为纯粹的启发式数字。这导致成本模型的更改往往是事后补救，滞后于代码生成的改进。经常在硬件或代码生成变更数周或数月后，才有人来调整启发式数字。

  3. 粒度低：许多后端的成本粒度很低。例如，add 指令的成本是 1。另一个指令可能比 add 稍贵，但并非两倍贵，然而成本只能是 1 或 2，缺乏中间值。有些后端（如 RISC-V）通过将所有成本乘以 100 来获得百分比粒度，但这导致了各目标自行其是的缩放方式。

  4. 简单的成本累加：许多地方只是简单地对 IR 指令成本求和，以此估算一个基本块的执行时间。这是一种非常简化的方式，可能无法准确反映实际硬件执行情况。

  5. IR 与硬件执行的差距：由于模型基于 IR 级别，它不一定能精确匹配硬件实际产生和执行的内容。例如，IR 可能对应多条机器指令，但后端通过融合等技术可能生成更简单的汇编代码，而简单的 IR 成本求和无法捕捉这种优化。

潜在的改进方向

认识到这些挑战后，我们思考了一些可能的改进方案。请注意，本节的目的不是提出最终解决方案，而是开启与社区的对话，探讨是否有更好的方法。

以下是我们考虑过的几种思路：

  • 提取 JSON 配置文件：创建一个简单的 JSON 配置文件，允许以更简单的方式覆盖后端的成本。这可以定义特定 IR 指令和类型的成本覆盖。这种方法便于通过机器学习或自动调优框架来大规模调整不同内核的成本。

  • 利用调度模型计算关键路径：与其简单累加基本块内所有指令的成本，不如利用 LLVM 中某些内核已有的调度模型，尝试计算给定块的关键路径长度。例如，在考虑循环向量化时，可以只累加关键路径上的成本，而不是所有指令的成本。公式上，这可以表示为：总成本 = Σ(关键路径上指令的成本)。

  • 基于代码生成进行更精确的成本推导：理论上，可以进一步深入代码生成阶段，理解后端最终会生成什么，并据此更精确地推导成本。这能更好地映射 IR 指令到真实的硬件效应（如指令融合）。然而，这种方法的主要问题是编译时间可能会显著增加。

社区讨论与总结

在演讲后的问答环节，社区成员提出了一些有价值的观点和建议：

  • 利用现有补丁：有社区成员提到 Simon Pilgrim 几年前的一个补丁，该补丁通过后端 lowering 并利用 TableGen 驱动的方法来生成成本表。这被认为是一种可行且有趣的方法，可能是改进的第一步。
  • 使用 MCA 进行估算：在讨论关键路径方法时，有建议指出既然已经进行了代码生成的繁重工作，可以直接使用 LLVM 的机器代码分析器（MCA）来估算成本。
  • 静态映射折中方案：作为完全代码生成和纯 IR 成本之间的折中，可以考虑建立一个从 LLVM IR 指令到潜在机器指令的静态映射，然后使用调度模型对这些机器指令进行成本估算。
  • 成本种类的扩展：目前成本模型主要使用一个启发式数字来代表一切（最初是吞吐量）。未来如果更紧密地结合调度模型，可能会更多地使用延迟、吞吐量等多种指标来指导成本计算。
  • 机器学习应用：成本模型的调优被认为是可能应用机器学习（ML）的领域，通过自动迭代来寻找更优的启发式参数或特征组合。

本节课中我们一起学习了 LLVM 指令成本模型的基础。我们了解到它是一个用于指导优化决策的快速启发式工具，具有可组合、可覆盖等优点。同时，我们也探讨了它在可维护性、粒度、精度以及与硬件匹配度方面面临的挑战。最后，我们介绍了几种潜在的改进思路，并看到了社区对此的积极讨论。改进成本模型是一个持续的过程，目标是在保持编译效率的同时，提供更准确、更易维护的决策依据，以生成更优的代码。

047：向量化性能问题与优化

概述

在本节课中，我们将要学习LLVM编译器在向量化过程中可能遇到的一个特定性能问题。我们将探讨向量化通常如何提升性能，但在某些特定场景下，例如处理跨步访问（stride access）时，向量化反而可能导致性能下降。我们将深入分析这个问题的根本原因，并了解针对不同硬件架构（如Arm的SVE和AArch64）的优化方案及其效果。

向量化会降低性能吗？🤔

向量化通常被期望用于提升程序性能。然而，我们发现，在某些跨步访问的场景下进行向量化，实际上会导致性能下降。

例如，在本幻灯片展示的测试案例中，向量化版本的性能实际上比标量版本更差。

我们发现，这种性能下降是由跨步访问的低效代码生成所导致的，并且我们已经为此问题提交了一个报告。

在本次讨论中，我们将探讨这个问题及其改进状态。虽然本次讨论主要关注针对SVE的AArch64架构，但我们相信这个问题也可能与其他架构相关。

当前向量化的问题：低效的地址计算 🔍

上一节我们介绍了向量化可能导致的性能问题，本节中我们来看看导致这个问题的核心原因之一。

当前由循环向量化生成的代码，其问题在于低效的地址计算。

对于SVE架构，跨步访问需要使用聚集（gather）指令。在当前向量化的代码中，操作的地址在循环内部使用向量指令进行更新，正如你在幻灯片示例中所见，向量寄存器Z1在每次迭代中都会被更新。

但是，如果我们能够识别出跨步访问的模式，就可以生成像这样更高效的指令序列。

以下是优化思路的代码描述：

// 低效方式（循环内更新向量地址）:
for (i=0; i<N; i+=VL) {
    addr_vec = base + stride * {i, i+1, ..., i+VL-1}; // 向量计算
    load(addr_vec);
}

// 高效方式（循环外初始化，循环内更新标量基址）:
addr_vec = initial_address; // 在循环外初始化向量寄存器
for (i=0; i<N; i+=VL) {
    load(addr_vec);
    addr_vec += stride * VL; // 使用标量指令更新基址
}

在幻灯片示例中，你可以看到基址寄存器X1在每次迭代中被更新。这段代码效率更高，因为它需要更少的资源用于地址计算。

问题改进状态与解决方案 🛠️

接下来，我将解释针对此问题的改进状态。

幸运的是，针对RISC-V架构（它拥有专用的跨步加载/存储指令）的补丁已经被提交。这个补丁在循环向量化过程中检测跨步访问模式，并将其分派给专用的跨步加载指令配方。

然而，AArch64架构并没有这些专用的跨步加载/存储指令，因此我们无法直接使用该指令配方。

所以，我们的贡献是创建了一个针对AArch64架构将配方合法化（legalize）的补丁。这个补丁用我们在上一张幻灯片中看到的指令序列替换了原有的配方。

实际上，对于引言中的测试案例，我们的方法相比当前的向量化方式，将性能提升了37%。

其他跨步访问向量化问题 📝

最后，我想简要提一下向量化跨步访问时遇到的其他一些问题。

以下是目前已知的其他相关问题：

  1. 具有可变跨步的循环未能被有效向量化。这与循环版本化有关，通常只有跨步为1的版本会被向量化，而其他情况则保持为标量代码。
  2. 在归纳变量简化器中，当将跨步访问的索引拓宽到64位时，可能会潜在使内存访问指令的数量翻倍。

我们已经将相关问题的编号列在了幻灯片上。

总结

本节课中我们一起学习了LLVM向量化中一个关于跨步访问的性能陷阱。我们了解到，尽管向量化旨在提升性能，但在处理特定内存访问模式时，如果地址计算不够高效，反而可能导致性能下降。我们分析了问题的根本原因，并探讨了针对AArch64架构的优化方案，该方案通过将循环内的向量地址计算转化为更高效的标量基址更新，成功提升了性能。最后，我们还简要了解了向量化跨步访问时面临的其他挑战。理解这些细微之处对于编写高性能代码和进行有效的编译器优化至关重要。

048：LLVM卸载基础设施教程

概述

在本节课中，我们将要学习LLVM卸载基础设施。这套工具集用于在外部加速器（如GPU）上运行程序。其核心思想是提供一个通用、可复用的框架，以便不同的语言前端和厂商编译器都能受益，避免重复实现。我们将通过一个快速导览，了解当前在Clang和LLVM中进行GPU编译的概况。

编译基础：将GPU视为嵌入式目标

上一节我们介绍了LLVM卸载基础设施的概览，本节中我们来看看如何将GPU编译视为一个标准的嵌入式目标工具链。

可以将GPU视为一个嵌入式目标。其工具链包含一个名为Clang的编译器。你向该工具链的编译器提供源文件，它会执行必要的步骤，生成一个可用的可执行文件。这种编译模式简化了与现有工具的集成。

以下是一个示例，它使用了通用的GPU内部函数头文件，使其在我们支持的所有目标之间基本通用。

#include <clang/gpu_intrin.h>
__global__ void kernel() {
    printf("Hello from thread %d on platform %s\n",
           __builtin_gpu_thread_id(),
           __builtin_gpu_get_platform());
}

你可以为每个GPU使用Clang工具链进行编译，例如针对AMD GCN或NVIDIA PTX，从而获得适用于所需硬件的可执行文件。目前对SPIR-V的支持尚不完善，主要缺少对非统一内存访问等功能的支持。

我提出这一点，是因为目前我正在基于这个思路推动LLVM基础设施来构建GPU运行时库。GPU确实有运行时库，例如数学库。在类似HIP的环境中，还有malloc和printf等。OpenMP也有一个相当复杂的运行时库。

构建这些库的方法就是使用直接编译的理念，将其视为一个标准的交叉编译工具链，就像你熟悉的任何其他嵌入式系统一样。

你可能也听过今天其他讲座提到的目标工具链库，例如multiarch。它的作用是将你的目标文件放入一个限定目录中，以便安装所有库而不用担心冲突。我一直在逐步将每个GPU运行时库迁移到这种方式，目前可能只缺libclc。

关于如何在CMake中构建，我们没有时间深入探讨，但卸载项目中有个缓存文件叫Offload.cmake，它就像是这方面的动态文档。如果你想了解具体做法，可以查看它。

异构编译与卸载工具

上一节我们介绍了如何将GPU作为独立目标进行编译，本节中我们来看看更常见的异构编译场景及其使用的卸载工具。

一旦我们构建了运行时库，就可以将它们用于编译更常见的GPU卸载语言，例如OpenMP、HIP、SYCL和CUDA。SYCL带星号是因为其上游功能尚不完整。

这里的主要区别在于这些是异构编译，它们将主机编译与一个或多个设备编译结合在一起。这给工具链带来了很多复杂性，为此我们有许多卸载工具来帮助管理。

你可以看到下面这个示例，它是之前示例的HIP版本。如果你加上-###标志，会看到很多步骤。我将快速概述这些步骤在做什么。

clang -### --target=amdgcn-amd-amdhsa -mcpu=gfx1030 -nogpulib example.hip

你提供源文件。如果使用-###，你会看到Clang前端生成了多个编译任务，产生一个主机对象文件和多个设备对象文件。每个Clang编译工具链都期望一个输出，但由于我们生成了多个设备对象文件，我们需要使用llvm-offload二进制工具将它们合并成一个。这只是一种包含元数据的二进制格式，用于在链接时重建原始的Clang任务。然后，我们将其嵌入到主机对象文件中，得到一个包含主机和设备代码的“胖二进制文件”。

接着，我们通过Clang链接器包装工具处理它，这个工具内部有很多魔法，但概括来说，它提取那些胖二进制对象，调用之前提到的直接目标编译来获得可执行文件，用注册代码包装它，然后将其合并到最终的a.out可执行文件中。这样，当你运行程序时，它会一次性初始化。

快速看一下嵌入的内容，它只是一个命名段。查看这个命名段，你会发现它只是一些元数据和一个二进制对象，并不太复杂。

如前所述，我们使用Clang链接器包装工具，这是一个为你完成所有工作的魔法步骤。但如果你愿意，也可以自己动手，假装自己是一个编译器驱动程序，使用llvm-offload二进制工具来提取文件。

以下是一个示例，展示了如何手动操作：使用llvm-offload提取文件，然后使用目标链编译器获得可执行文件，接着使用llvm-offload-wrapper将其包装在一个模块中，该模块包含类似hipRegisterModule的调用（具体名称我忘了）。这将被放入一个构造函数中，当用户启动程序时，该构造函数会被调用，然后你就可以访问你的GPU可执行文件了。

如果你不知道其中包含的架构，可以使用llvm-objdump配合卸载标志来查看。

卸载运行时库：通用GPU接口

上一节我们介绍了编译阶段，本节中我们来看看程序编译完成后，如何在GPU上实际运行它。

最近，LLVM项目增加了一个卸载运行时库，它是卸载项目的一部分，最初是从OpenMP支持中分叉出来的。

这个库旨在为所有我们打算支持的厂商GPU运行时提供一个通用接口。我说“通用”带星号，是因为我总是想警告人们不要试图创建真正通用的东西，因为这通常只会得到一个所有GPU功能的最小有用子集。因此，在需要的地方必须有扩展。目前我们支持AMD和NVIDIA的卸载，也有支持Intel GPU的补丁。这些主要是对实际厂商运行时的动态包装，但它提供了一个统一的接口，所以你只需要链接一个东西，它就能知道该怎么做。

这个库最初用于OpenMP，但我们发现其中大部分内容非常通用，实际上可以开始导出它。其他目标如SYCL也希望开始使用相同的接口。因此，最近有一个推动力，开始创建一个新库，将OpenMP拆分成自己的部分，并将通用部分提取到一个API中，这个API被称为liboffload。

那么，liboffload是什么？它是一个用于卸载到GPU程序的C API。我知道你们中有些人会想，我们不是已经有足够多的这类API了吗？答案可能是的，但在LLVM内部工作时，拥有一个我们可以控制的API是非常好的。

它仍在开发中，但可以说已经过了最小可行产品阶段，所以你肯定可以使用它，但有些东西可能会改变。像任何优秀的GPU运行时一样，它的每个API函数都以两个字符为前缀，比如VK、CL、CU。有人讨论过将其命名为LLVM卸载库，但我否决了，因为我不想在我的运行时代码中调用LO_。

像许多人之前做过的那样，我们通过TableGen生成头文件。你将API函数放在那里，给出描述以生成文档，然后它还会自动生成函数验证（例如检查参数是否为NULL等）和跟踪功能。当你运行程序时，可以选择性地获得打印输出，例如在哪个线程执行了什么、错误值等。这在调试时非常有用。

这对LLVM人员来说是老生常谈了。我们将这些东西安装在类似offload-api.h的头文件中，然后使用我们创建的新的TableGen可执行文件来生成实现。这将把它简化为如下形式：它进行跟踪检查，如果是，我们就发出函数名，调用实际运行时的实现，然后返回任何错误并打印出来。

从高层次看，许多GPU运行时看起来都一样，对象都暴露为不透明的句柄。如果你有一个设备，那只是一个你不知道的奇怪指针，你把它交给运行时，运行时为你做事。我们真的希望提供这个，让人们可以用于他们的编译器来发出卸载代码，对吧？

为了使这更容易，我们提供了丰富的错误消息，包含自定义的错误字符串，而不是像大多数运行时那样只给出一个通用错误和一个查找表。我们实际上为不正确的参数生成自定义的错误消息。

我们做的另一件事是，为每个函数提供一个接受代码位置的变体。因此，如果你是一个GPU编译器，你可以说错误发生在第52行，而不是在一个上万行的文件的某个地方。

在高层次上，我们基本上只是将所有东西都暴露为一个带有平台的设备，平台基本上是该设备的一个属性。因此，如果你有一个异构机器，上面有NVIDIA、AMD GPU甚至CPU，你只会看到一个大的设备列表，由你来挑选出你感兴趣的那个。

实践示例：加载与运行GPU内核

上一节我们介绍了通用的卸载运行时API，本节中我们通过一个具体的例子来看看如何使用它。

首先，我们需要一个可以在我们高级GPU上执行的GPU程序。我想使用之前的直接方法，这是严格的C代码。为了让它更高级一点，我将使用我编写的libc库运行时。

你可以看到这里我调用了GPU内部函数包含，这适用于几乎任何C风格的语言。我只是用#ifdef检查我们实际编译的目标。这是标准做法。

我声明了一个名为kernel的东西，它是一个GPU内核，将调用__builtin_gpu_printf，这会导致它以内核调用约定被发出到对象中，以便你可以在运行时查找并调用它。我们只是用线程ID打印那个平台。

现在，我可以使用之前的相同标志，为我的AMD GPU设置目标amdgcn-amd-amdhsa，并链接C库，然后得到一个可以在我的AMD GPU上运行的可执行文件。对NVIDIA GPU做同样的事情，得到一个可以在NVIDIA GPU上运行的可执行文件。

现在我们有了那个镜像，第一件事是编写一个可以加载它的程序。显然，我们只是从命令行读取那个镜像，然后我们想要初始化运行时。你可以看到我们调用了这个ol_init函数。

我们有一个很好的小助手宏来检查错误，这只会说，哦，如果有错误，就打印出来并退出。因为像许多其他运行时一样，这里的一切都会发出错误，你应该能够使用这种预检查机制。

假设你的初始化成功了，那么下一步是发现一个可以运行该镜像的设备。

由于所有设备都暴露为一个大的列表，你可以做的是遍历这些设备，然后找到一个可以加载你提供的二进制文件的设备。你使用ol_is_valid_binary函数和ol_iterate_devices函数来实现。这只需要一个回调函数和一些用户数据指针，你可以在回调函数中使用它来尝试加载。

这个镜像到这个设备上，如果成功，我们只返回false并退出循环，返回我们的设备；如果失败，我们继续下一个，直到找到一个。这样做的成本很低。

一旦我们知道了我们的设备，并且知道它会工作，我们需要做的就是获取一个程序和一个队列。你可以看到我们获取我们的设备，我们想要基本上加载那个镜像到我们的设备上。

这被称为ol_program_handle。所以我们调用ol_create_program，使用我们发现的、我们知道能够加载这个GPU镜像的设备，以及镜像来创建程序。这就是物理上将我们编译的二进制文件移动到设备上，以便我们可以在上面执行。

为了在GPU上做任何工作，你需要一个可以像FIFO风格推送任务的队列。所以我们继续创建一个队列。

最后，剩下的就是启动内核。因为我们知道这个内核的名字就是kernel，因为我这样命名它。我们需要做的是创建这个符号句柄，并在我们创建的程序中查找那个字符串名称。

这将遍历加载的二进制文件，找到这个内核符号的字面地址，并将其返回给我们，以便我们基本上可以调用它。

一旦我们有了那个，我们需要做的就是设置我们的参数和启动参数。在这个例子中，我只打算在四个线程上启动它（我本可以为更复杂的例子留出更多空间，但这只是打算在四个线程上启动它）。

我们调用ol_launch_kernel，将这个内核推送到这个设备上的这个队列，参数只是作为一个带有大小的结构体传递。一旦我们这样做了，我们需要做的就是同步队列，等待工作完成，然后清理我们之前创建的所有程序。所以我们销毁队列，清理程序，然后关闭设备。

现在完成了，我们需要做的就是编译它，并将我们编译的镜像传递给它。这只是一个标准的编译，链接到安装位置的卸载头文件和LLVM卸载库。

然后我们创建这个卸载二进制文件，并将我们的AMD GPU镜像传递给它，它成功地找到了我的AMD GPU，然后为我启动了这个程序。因为这是为了通用性而设计的，我实际上也有一台带有NVIDIA GPU的机器，我可以使用相同的程序，传递我的NVIDIA二进制文件，它会找到我机器上的那个NVIDIA GPU，然后使用相同的实现基本上执行相同的事情。

我认为这是一个非常有说服力的例子，展示了如今这一切是多么简单。你只需要一些源代码，将其编译成可以运行的东西。这只有几行C代码和不到50行的C++代码使用这个库，我基本上已经可以执行任何我想要的内核，将任务推送到那个队列，而且我不需要担心，你知道，所有这些都是通用的，我只有一个#ifdef，我只是放在那里以有一个不同的字符串。

未来工作与总结

上一节我们通过一个完整示例演示了基础设施的使用，本节中我们简要展望一下未来的工作方向。

是的，这基本上就是我要展示的全部内容。我可以更详细地介绍，但显然时间有限。所以我只简要谈谈我认为我们可以在这里做的一些未来工作。

SPIR-V正在兴起，但其中肯定有很多工作要做。我真的很想有一个SPIR-V版本的我的GPU内部函数头文件，因为那样我就可以轻松移植我所有的运行时库，这样你理论上就可以获得像我的printf和所有libc的东西，以及OpenMP，因为如果你没有看到那样的东西，它们几乎会自动发生，因为我已经将这些移植为通用的了。

SPIR-V还需要在运行时支持，我们实际上还没有即时编译它，但我不认为这会那么困难，只需要一些代码处理来调用一些外部工具。我不是MLIR专家，但我知道我们有一个像GPU方言的东西，可以做启动之类的事情，并且有后端实现会发出对HIP和CUDA的调用。我认为如果我们有一个可以卸载到这个的方言，那将非常酷。也许MLIR的人如果有兴趣可以联系我，但就像我说的，我不是MLIR专家。

显然，对于我们正在编写的实际API库，还有很多工作要做。你知道，我们需要更细粒度的初始化或功能选择。我真正想做的一件事是从卸载接口中提取更多libomptarget的东西。我们已经做了很多，但如果你查看源代码，仍然有很多东西上面写着OpenMP，我更希望它们都消失。另外，显然你希望这个库达到1.0版本并最终确定。我们需要更新我们的卸载文档，以便你实际上可以使用它。

但这基本上就是我要说的全部了。所以，如果你有兴趣使用这个或贡献代码，请联系我，我很乐意帮助你解决问题。

总结

在本节课中，我们一起学习了LLVM卸载基础设施。我们从将GPU视为嵌入式目标的基础编译开始，了解了如何构建运行时库。接着，我们探讨了异构编译的复杂性以及Clang和LLVM提供的卸载工具如何管理这些复杂性。然后，我们介绍了新的通用GPU运行时API——liboffload，它旨在为不同的厂商后端提供一个统一的接口。最后，我们通过一个完整的实践示例，演示了如何使用这套基础设施编译一个GPU内核，并编写一个主机程序来发现设备、加载二进制文件并启动内核。这套工具集使得在GPU上进行通用计算变得更加简单和高效。

049：新型编译器优化的自主发现

在本教程中，我们将学习如何利用大型语言模型（LLM）自主发现新颖的编译器优化。我们将探讨Google团队提出的Magellan项目，该项目使用AlphaEvolve框架，在LLVM编译器中自动化地寻找和生成高效的优化启发式算法。

概述

作为性能与编译器工程师，我们每天都会遇到许多优化问题，其中大部分是NP难问题。传统上，我们通过编写启发式算法来解决这些问题，但这个过程既费力又复杂。如今，我们可以利用人工智能，特别是大型语言模型，来自动化这一过程。Magellan项目正是将AlphaEvolve框架集成到LLVM编译流程中，以实现编译器优化的自主发现。

核心流程：AlphaEvolve框架

上一节我们介绍了利用AI进行编译器优化的动机，本节中我们来看看其核心实现框架AlphaEvolve。

AlphaEvolve是一个由Google DeepMind提出的通用代码生成智能体框架。其工作流程是一个迭代过程：

  1. 问题描述输入：首先，将待优化问题的文本描述作为提示词输入给预训练的大型语言模型。
  2. 程序生成：语言模型根据提示生成候选的优化程序（例如，一个编译器Pass）。
  3. 实际评估：生成的程序在真实硬件上被编译和运行，以评估其性能。
  4. 结果反馈：评估结果（如性能得分）和程序本身被存入程序数据库。
  5. 进化搜索：在下一轮迭代中，AlphaEvolve使用进化算法从程序数据库中筛选出表现最佳的程序，并将其上下文信息附加到新的提示词中，再次输入给语言模型。

经过多轮迭代，最终可以得到性能最优的程序。

集成到LLVM编译流程

了解了AlphaEvolve的基本原理后，我们来看它是如何具体集成到LLVM编译器中，形成一个自动化优化发现循环的。

以下是集成后的完整流程，每一步都紧密衔接：

  1. 新策略提议：AlphaEvolve在每次迭代中提出一个新的优化策略（程序）。
  2. 集成到LLVM：将该新策略（即一个编译器Pass）放入LLVM的lib/文件夹中。
  3. 构建新编译器：用集成了新Pass的LLVM源码构建出一个新的Clang编译器。
  4. 编译基准测试：使用新编译器编译一组基准测试程序，生成可执行文件。
  5. 性能评估：在真实硬件上运行这些可执行文件，进行性能剖析和评估。
  6. 反馈闭环：将性能得分作为奖励，连同编译日志和剖析结果一起反馈给AlphaEvolve。AlphaEvolve利用这些反馈和从程序数据库中采样的历史程序，为下一次迭代提出新的策略。

实验与应用案例

在将框架成功集成后，研究团队将其应用于多个具体的编译器优化问题。以下是几个关键实验及其结果。

案例一：函数内联优化（用于代码大小缩减）

这是MLGo项目先前用神经网络解决过的试点问题，被选作AlphaEvolve的首个基准测试。实验分为两种设置：

  • 设置A：基于预定义特征的局部启发式

    • 输入：与MLGo神经网络模型相同的预定义特征集。
    • 输出：一个布尔值，决定函数是否应该内联。
    • 结果：从一个简单的初始策略（拒绝所有内联）开始，经过两天迭代，生成的策略比上游LLVM的启发式算法实现了超过4%的代码大小缩减。

  • 设置B：基于任意LLVM API的完整启发式

    • 输入：只有LLVM中间表示（IR）。
    • 输出：AlphaEvolve需要实现完整的函数来决定内联与否。
    • 关键设计：将正确性检查与实际实现分离，确保生成的程序总是正确的。
    • 结果：虽然需要更多的试错，但在1.5天内达到了比设置A更好的结果。生成的代码已是可读的C++代码，长度仅为上游启发式代码的1/15，并能达到相似的优化效果。

此外，生成的策略展现了良好的泛化能力：

  • 时间泛化：在不同时间点的代码库上评估，性能依然稳健。
  • 应用泛化：将同一优化策略应用于其他内部基准测试，平均能实现超过8%的代码大小缩减，与之前的神经网络模型水平相当。

AlphaEvolve不仅复现了人工发现的特征，还自主发现了一些新的或简化了的特征，例如基于指令吞吐量的加权指令计数，以及对函数接口中指针类型转换的惩罚项。

案例二：其他LLVM性能优化问题

除了代码大小优化，团队还将该方法应用于更具挑战性的性能优化问题。

  • -ffast-math性能优化：在Clang的-ffast-math标志相关优化中，AlphaEvolve从一个最初导致性能倒退的启发式开始，最终自动找到了一个与人工编写启发式性能相当的策略。
  • 寄存器分配优先级队列：让AlphaEvolve为寄存器分配问题的优先级队列生成启发式。在搜索引擎应用上进行端到端评估后，它发现了一个简单的启发式，其性能与人工编写的版本持平。

案例三：扩展到MLIR优化问题

团队还将此方法从LLVM扩展到MLIR领域，评估了两个代表性问题：

  • 图重写（Graph Rewrite）：基于Google今年在MLIR会议上发表的EqualityGraph数据结构工作，让AlphaEvolve为其提出Yield指令的启发式。最终发现的启发式效率远高于人工版本。
  • 自动分片（Auto Sharding）：在Google举办的SPU 2025竞赛问题上进行评估。AlphaEvolve生成的解决方案在20支参赛队伍中取得了约第4名的成绩。

优势、挑战与未来方向

通过上述实验，我们总结了该方法的优势、面临的挑战以及未来的发展方向。

优势：

  • 提升生产力：自动化了费力的启发式算法设计过程，编译器工程师无需手动调整模型。
  • 采样效率高：通常只需数百次样本评估即可逼近最优启发式。
  • 生成可集成代码：直接生成人类可读、可修改、可调试的C++代码，易于集成到生产环境。

挑战：

  • 性能上限与收敛性：目前尚不清楚AlphaEvolve最终是否能收敛到全局最优解，以及能否在所有类型的编译器优化启发式上超越人类专家。

未来方向：

  • 尝试更多技术以突破性能边界。
  • 尝试解决“绿地问题”，即训练数据中未曾出现过的新颖优化问题。
  • 致力于基于开源大型语言模型和AlphaEvolve框架实现开源版本。

总结

本节课中我们一起学习了如何利用AlphaEvolve框架和大型语言模型实现编译器优化的自主发现。我们探讨了其核心流程、在LLVM中的集成方法，并通过函数内联、性能优化和MLIR优化等多个案例看到了其有效性和潜力。这种方法显著提升了优化探索的效率，并能生成可直接用于生产的代码，为编译器优化领域开辟了新的自动化途径。

050：数据布局中的不可能视角

在本节课中，我们将要学习LLVM中数据布局的描述方式，特别是如何处理那些无法通过标准对齐和大小规则来描述的复杂内存布局。我们将探讨当前方法的局限性，并介绍一种潜在的解决方案：引入显式的填充类型。

大家好，我是Justin Bogner，我在微软工作。我们目前正致力于将HLSL语言引入Clang和LLVM。今天我们将讨论数据布局，即我们如何描述数据在内存中的排布方式。

LLVM中用于此目的的主要工具是一个称为数据布局字符串或数据布局规范的东西。它基本上描述了各种类型的大小和对齐方式，例如指针大小等。通过这种方式描述的对齐约束，以及关于对象大小的隐式或显式规则，我们或多或少可以推算出数据将如何布局。

然而，有些情况无法通过这种方式解决。因此，我们必须在那些情况下使用显式填充。

例如，假设你有一个像 alignas 这样的关键字，结构体中的第二个 int 以某种奇怪的方式对齐。那么它如何被降低表示呢？我们只是在元素之间插入这种 i8 数组作为填充。我们最终还会在末尾添加填充，这使得这些结构体在容器等场景中表现良好。

现在，让我们稍微离题，谈谈HLSL中的构造，这也是我们正在做这些工作的动机。

HLSL有一个叫做常量缓冲区的东西。它源于许多旧的、遗留的设计思想，与早期GPU相关。但你可以将常量缓冲区想象为一组16字节的寄存器，附带一系列打包规则。对于整数、小向量和浮点数（标量和小向量），它们会尽可能紧密地打包在一起。

例如，这里有一个浮点数，后面跟着一个三元素浮点向量，它们会被打包到同一行或同一个寄存器中。然后这里有一个三元素浮点数，它在另一个寄存器中。接着是一个二元素向量，它放不下了，所以进入下一个寄存器。因此，这些数据会尽可能打包，但它们会尽量避免跨越这些寄存器边界。

另一方面，结构体总是从一个寄存器边界开始。所以即使它很小，即使后面有空间，也不会在它后面打包另一个结构体，它总是从一个寄存器边界开始。但我们仍然可以像之前一样，在它后面自由地打包标量。我们可以看到这里的 j 就紧跟在第二个结构体后面。

数组变得有点奇怪，因为在数组中，我们希望每个元素都独占自己的寄存器。所以它实际上会把数组拆开，在所有元素之间插入填充。但关键且让这些布局变得棘手的是，我们仍然可以在数组之后将其他数据打包到同一行中，所以那个 float 就放在里面。

因此，我们在元素之间有填充，但在数组之后没有填充。回想一下之前的 alignas 例子，它们是在元素之后添加填充，所以我们不能做完全相同的事情。但我们采用了相同的基本思路，即使用显式填充，只是我们必须使用一些非常扭曲的类型来实现。我不会展示这些类型的样子，它们非常丑陋。

现在，让我们谈谈如何识别那些显式填充。如果我们只有一个 i8 数组，这实际上可能非常模糊，因为你无法判断这到底是一个实际的字符串还是填充。在大多数CPU后端，这不太重要，因为我们根本无法访问它，可能没问题。但在DirectX和SPIR-V后端，我们实际上需要能够为元数据与驱动程序的通信等目的重现原始类型。在SPIR-V中，类型甚至编码了访问的逻辑索引和偏移量，因此我们需要能够获取这两者并跳过填充。

我们在这里所做的是，实际上创建了一个目标扩展类型，即显式填充类型。这是一个目标扩展类型，用于声明“这是填充，它这么大”。这在很大程度上是可行的，但有点尴尬，因为你需要为你支持的每个目标使用不同的目标扩展类型，但大体上可以应付。

但我认为，有理由在LLVM中引入一种更“一等公民”的显式填充类型。问题是，我们是否应该有一个 pad8 类型？这本质上将是 i8 的别名，但明确表示这是“填充”。这可以在几个地方使用：

  • 完全控制布局：例如，当你使用 alignas 或类似Vulkan属性 [[vk::offset]] 时，你可以完全控制数据最终的位置。
  • 聚合体的标量替换：在分解聚合体并提取其部分时，这些信息很有用，可以表明“这部分不重要”。对整个结构体的 memcpy 等同于复制每个成员，而你可以直接跳过填充部分。
  • 变换验证：在像Alive这样的工具中，当你试图判断两段IR是否在语义上完全等同时，拥有“这是填充，从中读取是poison”的信息会很有用。因此，我们可以直接跳过它，或者判断对整个结构的复制与对各个部分的复制是等价的。
  • 其他用途：在定义影子内存的消毒工具中，或者只是声明某个访问是非法的，也可能有其他用途。

我在这里想听听其他人的意见，关于你们认为这在其他哪些地方可能有用。这是我们应当推进的事情吗？我是否应该提交一个提案，将这个类型添加到LLVM？它是否具有普遍用途，能让我们清理目标布局相关的东西？还是我们应该只待在自己的沙盒里，保持现状也不错？

以上就是我今天的全部内容，非常感谢大家。

本节课总结

本节课我们一起学习了LLVM中描述数据布局的机制。我们首先回顾了标准的数据布局字符串如何描述类型大小和对齐。接着，我们探讨了HLSL常量缓冲区等复杂场景中，标准规则无法描述的布局问题，这需要通过显式填充来解决。然后，我们分析了当前使用目标扩展类型来表示填充的局限性。最后，我们讨论了一个潜在的通用解决方案：在LLVM中引入一个一等公民的显式填充类型（如 pad8），并探讨了其在布局控制、优化变换和验证等多个方面的潜在应用价值。

051：面向Python开发者的MLIR领域特定语言

概述

在本教程中，我们将学习PyDSL，这是一个基于Python的领域特定语言，它能够生成MLIR代码，并进一步编译为可在CPU或特定硬件上执行的目标文件。我们将介绍其核心概念、新功能以及开源进展。

背景介绍

上一节我们概述了PyDSL，本节我们来深入了解其基本工作原理。

PyDSL是一个基于Python的DSL，它通过生成MLIR来实现高性能计算。其核心机制是使用一个编译装饰器。

以下是其工作流程的核心代码描述：

@compile
def function_f(...):
    # Python函数体

装饰器会提取函数的Python抽象语法树，并将其与Python上下文结构一起传递给PyDSL的编译方法。该方法遍历AST并生成对应的MLIR代码，随后将MLIR进一步编译为目标文件。目前，PyDSL支持CPU和华为昇腾NPU目标。

新功能：模板装饰器

上一节我们介绍了基础编译流程，本节我们来看看一个类似于C++模板的新功能——内核特化。

为了实现内核特化，PyDSL引入了模板装饰器。它的工作方式与基础的@compile装饰器略有不同。

以下是模板装饰器的使用示例：

@template
def calc(...):
    # 函数体

@template装饰器仅提取函数的AST。实际的编译过程会延迟到函数调用时进行。在调用时，具体的模板参数会被注入到上下文结构中，然后传递给编译方法以生成特化的MLIR代码。

新功能：自动调优装饰器

除了模板，PyDSL还引入了自动调优功能，以帮助寻找最佳性能配置。

@autotune装饰器的工作方式与@template类似，它将配置参数注入到上下文结构中，然后调用编译方法生成特化的MLIR，编译为目标文件并运行，同时测量执行时间。它会穷举所有配置的笛卡尔积，因此可能耗时较长，仅建议用于调优较小的配置空间。

Triton互操作性支持

PyDSL现在支持与Triton的互操作，允许在PyDSL方法中调用Triton内核。

这是通过分析PyDSL方法中的调用点来实现的，目的是确定Triton方法的类型签名和计算表达式。然后，系统会获取Triton的AST，调用Triton编译器生成TTG IR，再通过适配器将Triton IR转换为上游MLIR的Linalg方言格式，并将其拼接到PyDSL的主模块中，最后一起进行编译。

功能改进与优化

在过去的一年中，PyDSL还引入了一些改进和便利功能。

以下是主要的改进点：

  • 张量索引：支持使用Python切片语法来实现内存子视图、张量提取切片和张量插入切片操作。
  • 内存布局：支持带步长的内存布局，以满足对数据布局有特定要求的后端。
  • 归约操作：引入了强大的归约操作，支持沿多个维度进行归约，并允许使用自定义的组合函数。

调用宏

自去年以来，PyDSL引入了调用宏，并在对NPU目标的支持中广泛使用。

调用宏本质上是将Python AST宏扩展为MLIR。它被设计为PyDSL的一个可扩展性功能，允许更灵活地生成底层代码。

开源进展

最后，整个PyDSL项目已过渡到开源开发模式。

以下是开源生态的组成部分：

  • 持续集成：建立了健壮的CI系统以支持活跃开发。提交PR时可以运行CI进行检查。
  • 问题追踪：所有问题和拉取请求都在开源平台上进行跟踪管理。
  • 社区交流：建立了Discord群组，欢迎所有希望参与PyDSL开发的人员加入。

总结

本节课我们一起学习了PyDSL的核心机制，包括其通过装饰器生成MLIR的流程、用于内核特化的模板功能、自动调优支持、与Triton的互操作性、过去一年的功能改进（如张量索引和归约操作）、调用宏的可扩展性设计，以及项目向开源模式的转型。PyDSL为Python开发者提供了一个强大的工具，以领域特定语言的方式生成高性能的底层代码。

052：理解linalg.pack和linalg.unpack

在本节课程中，我们将学习MLIR中linalg.pack和linalg.unpack操作的核心概念、工作原理及其在性能优化中的应用。我们将通过简单的例子解释它们如何通过数据重排来提升内存访问效率。

概述：为何需要Pack/Unpack操作？🚀

首先，我们通过一个例子来说明这些操作的价值。考虑一个非常著名的分块矩阵乘法算法。

在第一次迭代中，我们使用红色值来计算蓝色值的部分结果。在第二次迭代中，我们使用绿色值来获得完整结果。

观察其内存布局，情况看起来并不算太差。一些值彼此靠近，我们可以从中获得一定的局部性。缓存表现尚可。这是一个被广泛研究和熟知的算法。

然而，仔细观察，其中仍然存在一些不足。例如，红色值在内存中并非全部相邻。

打包布局的优势📦

在打包布局中，我们主动改变内存中的值排列。这样，我们就能始终进行连续的加载操作。

这是同一个矩阵乘法，但经过了一些重排。现在，红色值在内存中彼此相邻，我们可以高效地加载它们。同样，绿色值也彼此靠近。

在右侧的下一次迭代中，可以看到非常连续的浮点数。通过这样的算法，我们获得了极佳的缓存局部性。

这种布局被称为PE布局，而实现这种布局的算法通常被称为数据分块，正如我的同事Iga之前提到的。

Pack/Unpack操作的核心机制🔧

为了理解pack和unpack，需要将二维张量重新组织为四维张量。在我的思维模型中，我首先对张量应用一个视图。从一个逻辑上的二维张量开始，然后应用四维索引，就像你看到的红、绿、蓝、橙色块一样。

到目前为止，我们还没有改变任何内存布局，这只是改变了索引方式。

打包布局则获取这种索引方式，并将其实际体现在内存维度中。这里有一个简化的示意图，现在我们拥有四个维度而不是两个。当然，这些数据移动并非没有代价，我稍后会讨论如何分摊这些成本。

Pack操作示例📐

让我们看几个pack和unpack操作的例子。这里有一个pack操作。

它的一项功能是将分块大小填充到特定尺寸，例如，为了在GPU上使用张量核心。这是通过一个名为pad的参数来实现的。

这里，输入张量是3x1，而内部分块大小是5x2。该操作会在维度之间进行填充。可以看到，红、绿、蓝值现在并非彼此相邻，而是被分开了。

Unpack操作示例📤

既然我们可以创建这样一个填充后的张量，自然也有其逆操作，即将填充后的张量解包回原始状态。

这里，输入尺寸是1x1x2x3，输出尺寸是1x2。

内部高阶尺寸在某种程度上反映了输入尺寸。unpack操作会检查结果中的元素数量是否小于输入张量的元素数量，然后对输入的一个子集进行切片提取，而不是使用完整的维度。

动态形状支持🔄

这些操作也支持动态形状。这里是一个unpack操作示例。输入是动态的，输出也是动态的。

现在，内部分块尺寸有tile_h和tile_w参数。看起来像是属性的一部分，但实际上不是，这只是为了打印美观。它们是操作的参数。

降低此操作时，代码中会出现一系列张量维度操作。首先，创建一个空的张量。然后转置值，折叠形状。接着，因为需要能够进行填充的逆操作，必须获取输出张量的维度，然后按输出大小进行切片提取。

最后，使用linalg.copy将结果复制出来。

linalg.pack操作同样完全支持动态形状。

实现细节与观察🔍

在准备这些幻灯片时，我注意到一个有趣的现象：linalg.pack操作本身并不直接降低。但linalg.unpack可以。查看代码，会发现一条大约两年前的注释，指出insert_slice和extract_slice操作功能不够强大，不支持动态形状。但现在它们已经支持了。我认为只是目前还没有人需要直接降低pack操作。

单位维度的挑战⚡️

pack操作的一个主要痛点出现在处理单位维度时。单位维度是指大小为1的维度。在生成打包布局后，对外部维度应用分块时，最终外部维度经常会是1。

有几个模式会对单位维度进行特殊处理，因为它们经常出现在代码中。棘手之处在于，大小为1的维度是一个特殊情况，它不执行任何操作，但又必须考虑它。这里经常出现很多错误，不过我认为现在大部分问题都已解决。

观察这个操作本身，它实际上除了改变张量的视图外什么也没做，四个元素仍然存在，每个维度的大小都是1。然而，我们有一个需要在上游修复的测试用例或操作，并为此编写了测试。

现在，也可以在单位维度之间拥有非单位维度，并得到特殊处理。J在几周前实现了这个功能，非常方便。

性能收益📈

如果我们能利用数据布局并分摊内存移动的成本，就能获得显著的加速。需要说明的是，这是在Ega用于基准测试的Arm Neon指令集上测试的。

朴素的分块矩阵乘法已经相当快了。但当我们能够分摊数据移动的成本时，就能获得前所未有的速度提升。延迟降低范围从7.66%到惊人的92%。

有效使用指南🎯

如何有效使用这些操作？如果只有一个独立的计算内核，大多数情况下，切换到打包布局可能看不到性能提升。

要看到性能提升，需要转向像AI编译器这样的框架，能够将打包操作与生产者操作融合，从而隐藏与数据重排和内存移动相关的延迟。例如，我们使用IREE作为底层支持，它可以从中获得非常出色的性能。

未来方向与总结💡

目前还有哪些未探索的路径？回顾基准测试，大部分是计算密集型工作负载。我们也非常关注卷积，卷积的数据分块是我们目前尚未探索的领域，时间有限。

另一件事是关于linalg.pack和linalg.unpack操作的设计。它是基于属性的。如果看看其他项目，比如Triton，它们使用矩阵来表示这种内存重排。在我看来，那种方法比我们基于属性的方法优雅得多。如果我们探索类似的方法，或许可以自动将其应用于像linalg.generic这样的操作，然后用它来创建linalg.pack和linalg.unpack操作。这样，就可以用一个矩阵来表示你的内存层次结构，并自动将其应用于通用操作，这将非常棒。

本节课中，我们一起学习了linalg.pack和linalg.unpack操作如何通过改变数据布局来优化内存访问模式，从而提升计算性能。我们了解了它们的基本用法、对动态形状的支持、处理单位维度的挑战，以及在实际应用中分摊数据移动成本以获得性能收益的关键。

053：自动生成重写与降级模式 🧠

在本教程中，我们将探讨如何为 MLIR 中的不同方言自动生成重写和降级模式。MLIR 承诺提供共享抽象，允许你为特定领域编写编译器，并定义一次优化，即可在多个管道中复用。然而，现实是存在许多语义各异的算术方言，导致需要为每个方言重复实现优化，这既容易出错又耗费人力。我们将介绍一种利用计算能力而非人力的方法，通过合成工具自动生成这些模式。

问题背景：为何存在众多算术方言？

上一节我们提到了 MLIR 的共享抽象愿景。本节中，我们来看看为何在实践中，MLIR 生态系统中存在如此多不同的算术方言。

这些方言虽然都涉及算术运算，但在语义上存在重大差异：

  • 毒值语义：某些方言（如受 LLVM 影响的）引入了毒值概念，例如除以零或溢出可能产生任意结果，以便进行激进优化。
  • 多值逻辑：某些方言（如硬件相关方言）支持三值或九值逻辑。
  • 除法定义：对于除法的行为，有些是已定义的，有些则未定义。
  • 溢出处理：有些方言设计上不允许溢出，有些则可以。
  • 位宽成本：在某些上下文中（如硬件或 LLVM），增加整数的位宽成本可能很低或很高。

这些语义差异直接导致了优化规则的不同：

  • 结合律：大多数方言支持，但 arith 方言如果设置了 NSW（无符号溢出）标志，则可能不支持。
  • 乘幂转换：大多数方言可以将乘以 2 的幂转换为移位操作，但在 index 方言中，由于位宽未知，这可能产生毒值。
  • 乘除抵消：在不允许溢出的方言中，(x * c) / c 可以简化为 x；但在允许溢出的方言中，此优化不成立。

因此，为每个方言手动实现一套完整的优化（如指令合并）不仅工作量大，而且容易出错和遗漏。

解决方案愿景：自动合成模式

面对为众多方言手动编写优化规则的挑战，我们提出一个愿景：能否自动合成或生成每个方言的指令合并等价物？我们能否用计算能力替代人力？

为了实现这个愿景，我们构建并正在改进三款工具：

  1. 重写合成器：输入方言描述，自动生成一组重写规则。为了保证可行性，需要设置模式左侧和右侧操作的最大数量限制，并可能限制常量集（如仅 0 和 1）。
  2. 降级合成器：给定源方言和目标方言，自动生成降级规则。例如，将 shift_right 从 SMT 方言降级到 arith 方言时，需要处理移位量的边界检查，以避免引入毒值。
  3. 超级优化器：给定一个输入程序（使用特定方言）和该方言的描述，找出可以应用于该程序的所有重写规则。例如，将 x * 2 重写为 x << 1。

核心技术：枚举合成

我们使用枚举合成技术来实现上述工具。其核心思想很简单：给定一个规范（例如 2 * x），枚举所有不超过特定大小的程序，然后逐一测试它们是否在语义上等价。

一个可用的枚举合成系统需要三个核心组件：

  1. 枚举器：用于生成 MLIR 程序。
  2. 成本模型：用于筛选出比原始程序更优的候选程序。
  3. 等价性检查器：用于判断两个程序是否语义等价。

等价性检查器

我们使用 SMT 求解器（如 Z3）来构建等价性检查器。具体方法是，将两个 MLIR 程序编译到 SMT 方言，添加必要的粘合代码，然后交给 SMT 求解器判断它们对于所有输入是否产生相同输出。

通用枚举器

我们希望构建一个能枚举任意 MLIR 方言程序的通用枚举器，而不是为每个方言重写。我们利用 Guided Tree Search 库来实现。

以下是生成程序的基本思路：

  1. 为给定结果类型选择一个操作。
  2. 选择该操作所需操作数的数量。
  3. 为每个操作数递归地选择其类型，直到达到所需的操作数量。

为了确保生成合法的程序（例如，add 操作的操作数类型必须匹配），我们利用 MLIR 的 IR 方言来获取操作的元信息（操作数数量、类型约束等），从而指导枚举器生成有效的程序。

结合枚举器和等价性检查器，我们就可以执行枚举合成。这对于前两个工具（重写和降级合成器）是直接可用的。对于超级优化器，我们寻找与输入程序等价的程序（对于优化）或“精化”程序（对于降级，确保不引入未定义行为）。

核心算法：高效生成所有重写

我们承诺能生成特定大小范围内的所有重写规则。一个朴素的方法是枚举所有左侧程序，并为每个左侧程序枚举所有右侧程序。但对于 3 个操作的程序，约有 1000 万种可能，这意味着需要检查 100 万亿个候选对，这是不可行的。

关键在于，许多候选程序是冗余的。例如，x + 0 和 0 + x 都可以简化为 x。我们可以利用这种冗余来大幅缩减搜索空间。

以下是高效生成所有重写的算法步骤：

  1. 初始化：从大小为 1 的程序开始枚举。
  2. 分类：使用分类器找出所有等价的程序，形成等价类。
  3. 提取代表与可跳过模式：
    • 从每个等价类中选出一个代表（通常是最简单、操作数最少的程序）。
    • 将类中其他程序标记为可跳过，因为它们在任何上下文中都可以被其代表替换。
    • 这些“代表 -> 可跳过”的关系本身就是重写规则。
  4. 迭代枚举：在枚举更大尺寸的程序时，只使用当前尺寸的“代表”程序作为子表达式来构建新程序。同时，利用已积累的“可跳过”模式库，在生成时直接过滤掉包含这些低效子模式的程序。
  5. 使用 PDL 加速过滤：为了高效地检查一个程序是否包含任何“可跳过”模式，我们将这些模式编译到 MLIR 的 PDL 方言 中。PDL 可以将大量模式编译成一个高效的状态机，使得检查复杂度与最大模式大小呈线性关系，而非模式总数，从而极大提升速度。

通过这种迭代的、利用已知等价信息来引导后续枚举的方法，我们可以将搜索空间减少几个数量级，使得生成所有重写在计算上变得可行。

实践评估与结果

我们以 MLIR 中的 SMT 方言（约 30 个操作）为例进行实验，并允许使用常量 true, false, 0, 1。

以下是生成不同大小重写的结果：

  • 大小 1：共 372 个程序。我们生成了 100% 的重写，但只有 20% 的程序是“代表”（即最优形式），80% 的程序可被重写。同时提取出 115 个“可跳过”模式。整个过程约 1 秒。
  • 大小 2：利用大小 1 的“可跳过”模式，我们仅需枚举原计划 16% 的程序，过滤后仅需处理 8.9%。最终只有 4.5% 的程序是“代表”，并发现了约 1000 个新的“可跳过”模式。总耗时不到 1 分钟。
  • 大小 3：原始有 1400 万个程序。我们的算法仅需处理其中的 1.8%。最终 1.3% 为“代表”，并发现了约 43，000 个“可跳过”模式（即重写规则）。在 M2 MacBook 上耗时约 10 小时。

这些“可跳过”模式库也能加速超级优化和降级合成，通常有 3-5 倍的性能提升。

对于 arith 方言，由于其语义更复杂（涉及毒值），等价性检查更慢，生成时间比 SMT 方言长得多。

在降级合成方面，我们尝试为 SMT 到 arith 的 34 种操作（及不同类型）生成降级规则。大多数单操作降级可在 1 分钟内找到，双操作降级需要 2-10 分钟。我们还发现了一些需要 3 个甚至更多操作才能完成的降级案例。

总结与展望

本节课中，我们一起学习了如何利用枚举合成技术，为 MLIR 中的不同方言自动生成重写和降级模式。我们介绍了三款工具：重写合成器、降级合成器和超级优化器，并深入讲解了其核心算法——通过迭代分类和利用“代表/可跳过”模式来高效枚举所有可能的重写。

目前这些工具已能工作，但生成时间仍有优化空间。我们认为，对于简单的、重复性的重写和降级规则，不应再手动编写，而应使用合成工具自动生成。这是迈向在 MLIR 中大规模自动生成优化规则的第一步。

未来的工作包括：

  • 改进运行时性能。
  • 开发泛化算法，以便从具体的重写实例推导出更通用的规则。
  • 集成数据流分析，以支持更强大、上下文相关的重写规则。

通过自动化这些模式生成过程，我们可以提高编译器开发的效率与可靠性，并确保不同方言间优化规则的一致性。

问答环节要点

  • 测试生成：有观众询问该工具是否可用于生成测试用例。演讲者确认，他们已使用枚举器和检查器来测试 arith 方言，并发现了测试套件中的 4-5 个错误。方法是枚举所有大小不超过 2 的程序，并检查它们是否被正确优化。

054：基于MLIR的面向多面体引擎的LLM调度原语生成器

概述

在本教程中，我们将学习一个基于MLIR和LLM的自动化代码优化系统。该系统旨在利用大型语言模型为多面体编译器生成有效的调度策略，从而替代或辅助经验丰富的性能工程师。我们将首先回顾其基础——一个名为Pomorphous的基于MLIR的、由变换驱动的多面体编译器，然后深入探讨如何集成LLM来自动生成正确的MLIR变换方言代码。

1：背景与基础工作

本节我们将介绍本项工作的基础——Pomorphous编译器。理解其工作原理是理解后续LLM集成部分的关键。

这项工作是Robert Jin Man、Aan和Xin Yu的合作成果。本次演讲建立在先前一项名为Pomorphous的、基于MLIR的、由变换驱动的多面体编译器工作之上。

先前工作的高级理念是：由专家用户（例如性能工程师）向Pomorphous工具提供一个计算内核以及一些变换策略。该工具会在计算内核的上下文中验证这些变换，以检查是否存在任何依赖违规。如果用户提供的变换序列是合法的，工具将继续应用变换并相应地生成代码。如果变换序列是非法的，工具会尝试修复违规，同时尽可能保留用户的策略。

以下是一个使用PDL（Pomorphous的前端）编写的Polly基准测试andVT内核的简单示例。如果您昨天参加了我的演讲，就会知道PDL可以生成MLIR。

用户希望应用于内核源代码的变换策略在MVT_schedule方法中描述。在该方法中，首先匹配两个目标嵌套的for循环，用户希望先应用分块，然后融合，接着以32, 32进行分块，之后重新排序两个内层点循环，并并行化最外层的分块循环。

这个MVT_schedule展示了变换序列的可组合性及其与内核代码的分离。这是利用MLIR变换方言基础设施的主要优势。然而，一个关键区别是，这里的调度实际上并不转换有效载荷IR。相反，它会为后续的验证过程记录每个变换操作对应的变换矩阵。

整个代码库已通过PDS网站开源。

2：验证过程与多面体概念回顾

上一节我们介绍了Pomorphous的基本流程，本节中我们来看看其核心的验证过程。为了阐明Pomorphous如何最大限度地利用MLIR的现有基础设施，我们先快速回顾一些多面体概念。

多面体调度函数在语句级别表达。这里，语句S依赖于语句R。Cta(S)描述了执行时间戳，即每个S实例运行的时间，Cta(R)同理。

调度函数具有如下所示的格式，其中矩阵C中的系数允许我们在可能时通过倾斜或平移来纠正非法的变换。并且需要Farca Lambda来线性地构造约束。

那个复杂方程中的D描述了依赖多面体。

为了在需要纠正调度时找到正确的平移或倾斜值，我们通过利用MLIR中可用的Presburger集合求解器来求解方程组，以找到矩阵C中的系数。

3：利用MLIR构建依赖多面体

在初步了解之后，让我们深入探讨如何利用MLIR。本节将展示如何利用MLIR的分析基础设施轻松构建依赖多面体。

这张幻灯片演示了一个简单的内核，并利用操作1（对A[3]的加载）和操作2（对A[3]的存储）之间的读后写依赖关系，来展示我们如何轻松构建依赖多面体。

我们通过编写一小段代码来调用MLIR中针对这对内存访问操作的checkMemrefAccessDependence方法，从而利用MLIR的分析基础设施。

我们需要为Uar循环（层级0）、内层arc5循环（层级1）以及最后的公共循环层级（+1，即层级2）调用此方法三次。因此，构建了两个依赖多面体：一个在层级0，另一个在层级2。这得益于MLIR，得以精确且准确地完成。

4：转换为求解器输入格式

在调用MLIR的单纯形求解器之前，我们编写了一个compute_farcus_rhs方法。该方法将上一张幻灯片中看到的依赖多面体转换为二维整数矩阵表示，如最右侧所示。

矩阵的行代表循环迭代器和程序参数，列是farcus乘数。类似地，使用fars_leftenci计算，我们寻找的矩阵C的系数是二维整数矩阵的列，并且我们寻求字典序最小解。

将变换矩阵和依赖多面体转换为调用单纯形求解器之前的适当整数矩阵布局的代码大约有850行。总而言之，我们以极小的努力就能够构建一个完整的多面体验证器，再次感谢MLIR。

5：引入LLM生成调度策略

在初步成功后，我们的领导团队要求我们探索是否可以用LLM来为Pomorphous生成优化策略，以替代经验丰富的用户。此外，我们还希望了解LLM是否能生成语法正确的MLIR代码，以替代PDL这样的DSL。

因此，这是演讲的第二部分：一个生成式MLIR LLM规划与变换IR生成器。

这是这部分工作的高级设计。它包含一个路由器、几个代码生成代理，以及作为守护最终正确性大门的验证器Pomorphous。相当简单。再次强调，该系统的目标是为Pomorphous生成变换IR。现在，让我们首先深入了解路由器。

6：路由器设计

路由器本质上是一个检索增强生成数据库。它接收待优化的输入内核，将其分类为与其数据库中前K个参考程序最相似的程序。数据库中参考程序的优化策略由人类专家提供。

分类网络是一个简单的、经过预训练的CodeBERT MLP Softmax架构，使用LLM生成的代码和真实数据进行训练。

路由器或LLM将从数据库中获取关于如何优化某些参考程序的确定性知识，并将这些知识应用于优化当前的输入代码。这是LLM的非确定性部分。

路由器会生成关于如何优化输入内核的人类可读策略（如图所示），并且它也能沿着热图中的对角线相当好地对程序进行分类。

7：LLM生成MLIR代码与实验结果

路由器与协作的代码生成代理共同输出语法正确的MLIR代码。以下是生成的MLIR变换方言代码。

一些实验结果。是的，我翻页太快了，但无论如何，早期的实验结果表明，所有30个Polly基准测试都生成了代码并正确执行。这对LLM来说是一个鼓舞人心的好成绩。其中27个程序的运行时性能优于或等于仅使用GCC13 -O3编译的代码，14个程序的运行时性能优于或等于最初在Pomorphous论文中发布的人工编写的调度方案。

总结

本节课中我们一起学习了如何将基于MLIR的多面体编译器Pomorphous与大型语言模型相结合，构建一个自动化代码优化系统。我们回顾了Pomorphous如何利用MLIR基础设施进行依赖验证和调度纠正，并深入探讨了LLM如何通过检索增强生成和协作代理来生成有效的、语法正确的MLIR变换策略。初步实验表明，该系统能够为基准测试生成正确且往往更优的代码，展示了AI辅助高性能代码生成的潜力。

055：Windows平台上的LLVM工具链中缺失的关键功能

概述

在本节课中，我们将回顾LLVM工具链在Windows on Arm平台上的进展，并重点指出当前存在的一些关键功能差距。我们将从已实现的功能开始，然后深入探讨在调试器、链接器、编译器和其他工具方面仍需完善的地方。

已实现的功能

上一节我们介绍了课程的目标，本节中我们来看看LLVM工具链在Windows on Arm上已经取得的成就。

自LLVM 16版本以来，我们已经拥有了原生的LLVM工具链。其核心组件均能正常构建和运行，代码生成质量、运行时性能表现良好，并且获得了广泛采用。目前的目标是弥合最后的差距，使其在功能上与AArch64 Linux以及Windows x64的开发体验看齐。

以下是当前可用的主要组件列表：

  • Clang/LLD：能够生成原生的Arm64位二进制文件。
  • Clang-cl：提供与MSVC的命令行兼容性，以及整体的ABI兼容性。
  • Flang：Windows on Arm平台上的原生Fortran编译器，这是LLVM工具链的一个亮点，因为它是该平台上的首个Fortran编译器。
  • LLDB调试器：支持原生调试工作流。
  • OpenMP：已启用支持。
  • MinGW工具链中的LLVM：在GCC工具链尚未就绪的情况下，MinGW工具链目前使用LLVM作为支持。

LLVM成为生态核心的原因

那么，LLVM是如何成为Windows on Arm生态系统支柱的呢？这主要得益于其共享的后端和前端设计。

  • 共享后端：Arm64后端在Windows、Linux和macOS之间共享。针对x64平台已实现的代码视图（CodeView）和结构化异常处理（SEH）支持，大部分可直接复用。Windows on Arm主要需要处理大量ABI相关的工作，在修复一些问题后便已就绪。
  • 共享前端：C++、Rust、Fortran、MLIR等语言的前端支持也是共享的。这使得这些语言的编译器能够为Windows on Arm平台提供支持，其中一些虽处于实验阶段，但如C、C++、Fortran、Rust等已在该平台上可用。

这种共享架构带来了显著的性能优势，并且在Windows on Arm上尤为重要，因为MSVC工具链缺少一些原生LLVM工具链具备的功能。

生态系统采用情况

接下来，我们看看LLVM在Windows on Arm生态系统中的实际应用。许多重要项目都在使用它。

例如，Python这一主要项目使用了Clang/LLVM库。同样，Chrome、Edge、Firefox等团队也广泛使用Windows on Arm编译器来为这些项目提供原生的Arm64 Windows二进制文件。

我们LLVM团队一直密切关注Windows on Arm生态系统，并建立了一个名为“Works on Windows on Arm”的网站，用于追踪不断增长的原生应用。如果你想了解哪些应用已原生支持该平台，可以参考这个网站。

性能表现

在性能方面，LLVM的表现非常出色。它不仅超越了模拟运行的性能（这是显而易见的），在SPEC2017基准测试以及各种实际工作负载（如编译时间对比）中也优于MSVC。

一个突出性能优势的项目是Blender，使用LLVM（Clang）相比MSVC带来了大约20-30%的性能提升。

至于Flang的性能，由于它是目前唯一的Fortran编译器，我们尚未在Windows on Arm上使用SPEC2017进行验证，但目前还没有用户提出相关投诉。

然而，需要指出的是，Windows平台（包括Arm）的整体开发体验速度相比Linux和macOS较慢。构建LLVM等项目时，NTFS文件系统开销或Windows进程创建开销会导致速度下降，对于需要生成大量目标文件的项目尤其如此。我们需要寻找新颖的方法来解决这些问题，例如探讨将Clang作为服务器使用的可能性，但这些想法尚未深入讨论。

关键功能差距：调试器 (LLDB)

现在，让我们进入核心部分，探讨需要投入时间解决的实际差距。我们发现了LLDB调试器存在最多的问题。

以下是已识别出的部分关键差距列表：

  • Arm64EC不支持：Arm64EC是微软开发的ABI，允许在同一二进制文件中混合x64模拟代码和Arm64原生代码，并可以来回跳转。LLDB目前不支持这种混合模式调试。此外，如果一个二进制文件使用了x64模拟，原生LLDB目前也无法调试它。
  • Armv9功能缺失：虽然目前市面上还没有搭载Armv9的Windows on Arm笔记本电脑，但云服务器硬件和QEMU模拟器已支持所有Armv9功能。LLDB目前缺少SVE/SME寄存器可见性支持。
  • 硬件观察点限制：硬件观察点虽受支持，但由于Windows内核的限制，目前只能设置一个。我们已向微软报告此问题，他们正在尝试修复。同样，硬件断点也存在来自Windows内核的限制，需要底层修复后才能在LLDB中实现。
  • DAP体验不稳定：DAP是LLDB的一个组件，允许其与VS Code等IDE通信。目前上游LLDB中的DAP体验非常不稳定，我们因此禁用了许多相关测试。可能存在一些下游公司的补丁，但上游版本问题较多。
  • 功能对比缺失：与Linux或macOS版本相比，LLDB在Windows上缺少许多功能。
  • PDB调试体验：LLDB需要设置PDB读取器插件（微软PDB读取器或原生PDB读取器）。虽然近期有一个PR修复了使用原生PDB读取器时失败的测试，但整体调试体验仍有差距。例如，使用MSVC可以进行“编辑并继续”，但使用LLDB（即使配合VS Code）则无法实现。此外，Clang生成的PDB调试信息有时会省略变量或将它们标记为已优化，这在用LLDB或MSVC调试器调试时都会出现问题。我们需要在Windows平台（特别是Windows on Arm）上对PDB调试体验进行大量测试和检查。

关键功能差距：其他工具

除了调试器，其他工具也存在一些差距。

LLD链接器：

  • 与LLDB类似，LLD也不支持链接Arm64EC代码。这意味着你可以使用Clang编译Arm64EC代码，但无法用LLD进行链接，必须使用微软链接器。原因是链接Arm64EC二进制文件所需的某些信息未公开，这限制了我们实现LLD支持。我们一直在与微软沟通，但不确定相关信息是否会很快公开。
  • LTO（链接时优化）：我们尚未充分验证在Windows on Arm平台上使用LLD进行链接时优化的能力。如果有人有相关信息或发现缺失环节，请与我们分享。

Flang编译器：

  • Flang为Windows on Arm提供了Fortran编译器，但仍存在一些障碍。例如，它不支持类似MSVC风格的命令行驱动程序（像clang-cl为C/C++所做的那样）。这意味着在混合使用Flang和MSVC链接器或需要复杂构建系统的项目中，集成Flang会面临问题，这也是我们未能在Windows上成功运行Flang的SPEC2017基准测试的原因之一。我们需要找到方法为Flang启用Windows上的命令行兼容性。

其他：

  • Sanitizers（消毒剂）：目前不支持Sanitizers（如ASan、UBSan等）。我们知道这个问题已有几年，但每当开始着手修复时，优先级总会发生变化。我们希望在未来支持该功能的LLVM版本中将其加入。

未来工作：Armv9启用

展望未来，Armv9的启用是重点。我们知道在未来6-18个月内，将有更多搭载Armv8.3+或Armv9功能（如SME、MTE）的硬件出现在Windows平台。LLVM已经支持大部分这些功能（主要在Arm Linux上）。我们需要等待Windows内核提供支持以及硬件上市，然后进行测试并修复剩余问题。

测试、CI与发布

在测试和持续集成方面，我们运行着多个构建机器人（bot）：

  • 用于测试Flang的单阶段构建bot。
  • 用于测试Flang/OpenMP编译器组合的bot。
  • LLDB测试bot。
  • 一个不太稳定的两阶段构建bot（运行在Surface Pro平板电脑上，硬件不太适合此类测试）。

最近，我们成功在Windows on Arm上启用了LNT和LLVM测试套件，并有一个暂存bot在运行相关测试。我们计划与合作伙伴共同努力，修复目前尚不支持的其他测试（如GFortran测试、单源/多源测试）。

关于发布，我们自LLVM 16起就开始制作Windows on Arm版本，并使用上游脚本。我们改进了脚本以支持该平台，最近还添加了基于PGO的版本。几个月前，GitHub Actions提供了原生的Windows on Arm运行器，我们计划很快将发布流程迁移过去。

至于SVE测试和验证，我们拥有支持Armv9的Azure云硬件，但资源有限，尚未能测试SVE验证工具是否能在该云基础设施上工作。

生态系统增长：GCC支持

生态系统增长对我们至关重要。最近，GCC实现了对Windows on Arm的支持（尽管仍处于早期交叉编译阶段，C++支持因结构化异常处理补丁仍在进行中而缺失）。这解锁了构建Git等依赖GCC的应用的能力，并有望帮助我们运行多源/单源测试以及GFortran测试。

总结

本节课中，我们一起学习了LLVM工具链在Windows on Arm平台上的现状。我们回顾了已取得的显著进展和广泛采用，也深入探讨了在LLDB调试器（如Arm64EC、Armv9支持）、LLD链接器、Flang编译器以及Sanitizers等方面存在的关键功能差距。此外，我们还了解了未来的工作重点（如Armv9启用）、当前的测试CI基础设施以及生态系统的最新发展（如GCC支持）。解决这些差距需要社区和合作伙伴的共同努力，以进一步完善Windows on Arm平台上的LLVM开发体验。

056：理解MLIR崩溃再现器

在本节课中，我们将要学习MLIR中的崩溃再现器。这是一个强大的调试工具，能帮助我们在复杂的编译流水线中定位和复现错误。我们将从基本概念开始，逐步深入到如何集成和使用它，最后探讨其内部工作原理以确保我们的方言能与之良好协作。

调试与复现错误的基础

上一节我们介绍了课程概述，本节中我们来看看调试任何程序或复现任何错误都需要三个关键要素。

以下是三个关键要素：

  1. 从有问题的程序中得到的实际输出，即程序给出的错误答案。
  2. 你期望看到的正确输出。
  3. 一种可靠的方法来证明你得到的输出与期望的输出不符。

这个理念并不新鲜。例如，在C++单元测试中，断言的失败会告诉你哪里出错了。在MLIR的Lit测试中，FileCheck会指出转换后的输出与预期不匹配。你只需运行测试来证明不匹配。

然而，当面对一个完整的编译流水线时，情况就不那么理想了。通常，用户会报告编译失败。此时，复现用户的错误变得困难，因为可能涉及代码分享、环境差异和版本问题，导致巨大的沟通开销和复现障碍。

崩溃再现器简介

上一节我们讨论了调试完整流水线的困难，本节中我们来看看MLIR如何通过崩溃再现器来简化这一过程。

让我们通过一个实际例子来理解。假设convert-vector-to-llvm这个Pass中存在一个bug，该bug会影响向量点积操作。如果我们运行这个Pass，可能会得到一个错误，但错误信息可能不够清晰，难以定位问题。

在这种情况下，我们可以让MLIR生成一个崩溃再现器。通过向mlir-opt传递额外的--pass-pipeline-crash-reproducer参数并指定一个文件名，MLIR会在发生错误时生成一个.mlir文件。

运行后，我们得到了与之前相同的错误输出，但同时MLIR提示生成了一个repro.mlir文件。让我们看看这个文件的内容。

生成的再现器文件大致分为两部分。第一部分包含了复现刚才失败所需的IR。第二部分是一个类似JSON的结构，被称为“外部资源”。外部资源是一种可以附加到IR上的元数据，而无需内联在IR本身中。一个常见的用例是存储模型权重等二进制数据，避免每次序列化都内联它们。

在这个外部资源中，pipeline_builtin字段记录了导致错误的Pass流水线（例如convert-vector-to-llvm）及其选项。下方还记录了运行mlir-opt时的Pass管理器和MLIR上下文选项。

一旦我们有了这个再现器文件，就可以用它来复现错误。只需运行mlir-opt --pass-pipeline-crash-reproducer=repro.mlir，就能得到完全相同的错误。关键在于，我们不再需要提供原始的Pass流水线或触发错误的原始IR。这个崩溃再现器是一个完全自包含的复现包，包含了重现错误所需的一切。

两种类型的再现器

上一节我们介绍了崩溃再现器的基本用法，本节中我们来看看MLIR生成的两种不同类型的再现器：完整再现器和局部再现器。

完整再现器包含初始IR以及导致失败点之前运行的所有Pass。局部再现器则只包含失败发生前一刻的IR以及触发失败的那个（或那几个）Pass。

为了演示区别，我们运行一个更复杂的流水线，其中包含已知会触发bug的convert-vector-to-llvm Pass，但前后还穿插了其他Pass（如convert-vector-to-scf和cse）。

如果生成完整再现器，其IR与我们最初启动流水线时的IR几乎相同，Pass流水线也完整记录了所有运行的Pass。如果生成局部再现器，其IR会有所不同，因为它反映了在失败点之前、其他Pass（如convert-vector-to-scf）已经施加了转换之后的状态。同时，其Pass流水线只包含导致失败的那个Pass（convert-vector-to-llvm）。

局部再现器可以快速缩小导致失败的Pass范围。但并非所有情况都适合使用局部再现器，我们稍后会解释原因。

如何集成到编译器中

上一节我们了解了两种再现器，本节中我们来看看如何将崩溃再现器功能集成到你自己的编译器中。

集成过程非常简单，MLIR提供了清晰的API。你需要使用Pass管理器，并调用enableCrashReproducerGeneration方法。你需要传递输出文件路径和一个布尔标志，该标志指示是生成局部再现器还是完整再现器。这完全对应于mlir-opt的命令行选项。

此外，你还可以让MLIR在任何情况下（即使没有失败）都生成再现器，但这需要你明确提供要运行的Pass流水线和要运行的操作。

接下来，我们看一个在真实编译器中良好集成的例子：Triton。Triton通过环境变量TRITON_REPRODUCER_PATH来暴露此功能。如果用户设置了这个变量，编译器会立即启用崩溃再现器生成，并将输出保存到指定路径，同时设置为生成局部再现器。

一个细节是，生成局部再现器时，需要在MLIR上下文中禁用多线程。即使没有提供环境变量，Triton的Pass管理器也会启用崩溃再现器生成，但会将输出导向标准错误流，这样如果发生崩溃，再现器信息会打印到控制台，而不会写入文件。重要的是，必须在运行任何Pass之前进行这些设置。

使用再现器的建议与技巧

上一节我们学习了如何集成，本节中我们来看看一些使用崩溃再现器进行高效调试的实用建议。

尽管崩溃再现器是自包含的，但它们可能包含许多与复现错误无关的信息。为了获得最精简的复现，首先应该尝试减小问题规模。这适用于任何调试场景。

以下是减小问题规模的步骤：

  1. 从Pass流水线中移除已知无影响的Pass。
  2. 更谨慎地移除那些“安全”的Pass，如规范化或公共子表达式消除。这些Pass通常使IR更简洁，但你不应依赖它们来保证IR的正确性。不过需要注意，如果你的某个Pass模式匹配依赖于规范化后的IR，则需要保留它。
  3. 开始手动修剪IR中的操作和值，同时确保错误仍然出现。目标是使IR连接更小。

如果经过上述步骤后问题仍不清晰，建议使用完整再现器而非局部再现器。因为局部再现器只捕获失败前一刻的IR状态，但问题可能早在流水线前期就被引入，直到后期才触发失败。使用完整再现器可以获得流水线的全貌。

使用MLIR-Reduce自动化精简

上一节我们提到了手动精简IR，本节中我们来看看如何利用mlir-reduce工具来自动化这一过程。

mlir-reduce是一个用于精简包含错误的输入IR的工具。它的工作原理是尝试应用各种缩减策略来使IR变小，并在每次缩减后检查错误是否仍然存在。

要使用mlir-reduce，你需要提供一个“兴趣脚本”。这个脚本的作用是告诉mlir-reduce，在应用了某种缩减后，当前的IR是否仍然“有趣”（即是否还包含你要找的bug）。

以下是一个简单的兴趣脚本示例：

#!/bin/bash
mlir-opt $1 --pass-pipeline="builtin.module(convert-vector-to-llvm)" 2>&1 | grep "vector.contract"
exit ${PIPESTATUS[0]}

这个脚本运行指定的Pass流水线，并检查输出中是否包含vector.contract错误信息。如果找到（即grep成功，脚本因PIPESTATUS[0]非零而退出码为1），则mlir-reduce认为当前IR仍“有趣”，会继续尝试缩减。

你可以这样调用mlir-reduce：

mlir-reduce --test=interesting_script.sh repro.mlir --traversal=topdown

兴趣脚本的定义完全由你决定。你甚至可以不使用mlir-opt，而是编译并运行IR，然后检查数值错误等。关键是脚本必须以退出码1表示“有趣”（bug仍在），以退出码0表示“无趣”（bug已消失）。

需要注意的是，mlir-reduce主要精简IR本身，而不会自动精简Pass流水线。流水线的精简通常需要手动完成。

确保方言与再现器兼容

上一节我们探讨了使用工具，本节中我们来看看崩溃再现器的内部工作机制，这对于确保你自定义的方言能与之良好协作至关重要。

当启用崩溃再现器时，Pass管理器的工作方式有所不同。它会在一个独立的线程（称为崩溃恢复线程）中运行Pass。这样，如果运行Pass时发生严重错误，该线程可以通过长跳转回到运行Pass之前的状态，而不会导致整个应用程序崩溃。

这种机制对你的方言实现有两个主要影响：

首先，确保你的Pass选项是可打印且可序列化到流的。这似乎显而易见，但容易被忽略。例如，如果你将Pass选项打包在一个C++结构体中，打印流水线到再现器时，打印机可能不知道如何打印这个结构体，从而导致令人困惑的崩溃。因此，建议将结构体选项展开为独立的基元类型选项。

其次，需要小心处理在Pass运行期间或方言内部创建的任何线程局部状态。因为线程局部状态对象是在崩溃恢复线程的生命周期内创建的，一旦该线程结束，再次访问该对象会导致释放后使用错误。一个具体的例子是DistinctAttr存储分配器曾经使用线程局部存储来避免竞争，这已得到修复。

总结与后续步骤

本节课中我们一起学习了MLIR崩溃再现器的强大功能。我们从终端到端编译器失败这一常见且令人沮丧的场景开始，看到了如何使用崩溃再现器来更可靠地复现这些失败，并开始迭代以修复错误。

我们演示了如何将其集成到真实的编译器中，特别是Pass管理器上需要调用的方法，并以Triton为例展示了如何通过杠杆和标志向用户暴露此功能。

我们讨论了使用再现器的建议，如何精简问题规模，以及如何使用mlir-reduce。最后，我们探讨了为了确保方言实现与崩溃再现生成良好协作而必须了解的内部机制。

接下来你可以：

  1. 尝试使用再现器并将其集成到你的编译器中。可以参考Triton的做法。
  2. 鼓励用户在提交错误报告时附上再现器。这可以大大简化调试过程。
  3. 尝试添加一些Lit测试来练习再现器，特别是当你做一些可能与之交互不佳的事情时（如使用线程局部存储或以奇怪方式指定选项）。更广泛的测试有助于使该功能更加健壮。

如果你想了解更多关于调试MLIR的工具和技巧，强烈推荐Christopher的演讲，其中包含了许多实用且经过验证的建议。

057：LLVM对高效沙盒化的支持

概述

在本节课中，我们将要学习一种名为“轻量级故障隔离”的新特性。这是一种用于进程内隔离的新功能，其RFC已获批准，我们正致力于将其引入LLVM。这项技术旨在为那些我们对其内存安全性存疑的代码（无论是第三方还是第一方原生代码）提供高效的进程内隔离。

引言：内存安全的挑战与现有方案

上一节我们介绍了LFI的概述，本节中我们来看看它要解决的核心问题。如今，内存安全是一个公认的重大挑战，我们需要更好的工具来应对它。不幸的是，目前我们处理此问题的选择非常有限。

以下是当前主要的几种方案：

  • 不健全的缓解措施：例如栈金丝雀、控制流完整性等。这些措施的优点在于它们能与现有代码协同工作，只需启用编译器标志即可。但问题在于它们经常被绕过，并且很难评估其实际收益。用户付出了性能代价，但安全收益不明确。
  • 用安全语言重写代码：这能提供健全的安全性，但通常伴随着巨大的工程成本。你需要重新构建一个系统，并且这可能与开源社区的协作模式不兼容。
  • 基于进程的沙盒化：这能提供健全的安全性，并且适用于现有代码。但其缺点在于性能开销巨大（如上下文切换、创建开销），难以扩展，并且破坏了开发者的原有体验，将单一应用变成了分布式系统。

轻量级故障隔离的愿景

上一节我们分析了现有方案的不足，本节中我们来看看LFI如何提供一种新的选择。我们期望找到一种兼具两种优点的方案：既能提供健全的安全性，又能与现有代码协同工作。

LFI的灵感来源于在Firefox中使用WebAssembly对第三方库进行沙盒化的实践。WebAssembly提供了一种高效的隔离模型，其上下文切换速度极快，启动时间短，将隔离视为一种语言和编译器构造，而非操作系统层面的介入。该方案自2021年起已在Firefox中部署。

然而，WebAssembly有其设计上的局限性，这限制了其对现有代码的兼容性和性能表现。因此，我们开始探索一种更低层次的沙盒化技术。LFI正是在此背景下诞生的，它旨在提供高性能、高兼容性、强安全性和简单性的低层次沙盒化方案。

我们与Android团队合作，成功地将一个解码器库用LFI进行了沙盒化，性能开销仅为4%。这证明了LFI的潜力。LFI不是内存安全的唯一答案，但它是一个强大的新工具，可以与边界安全等其他LLVM能力相辅相成。我们计划在2026年将其引入Android系统。

LFI技术架构：以AArch64为例

上一节我们了解了LFI的愿景和背景，本节中我们来看看它的具体技术实现，首先以AArch64架构为例。

LFI的整体编译器架构涉及创建一个新的目标子架构（例如AArch64的子集）。在编译管道的最后阶段（汇编器/MC指令层面），会进行一些模式重写，使生成的指令符合一个安全的架构子集。这种方法的好处是能处理手写汇编和内联汇编。

编译过程只需保留少量寄存器并添加重写阶段。编译后，你会得到一个经过重写的ELF二进制文件。你可以运行一个静态分析验证工具来确保所有指令都符合安全子集。然后，该二进制文件将在LFI运行时环境中加载和执行。

执行环境被限定在一个4GB的虚拟地址空间区域内。该区域分为不可写的代码段和不可执行的数据段。限制访问的核心思想是：将一个基地址存储在专用寄存器中，并确保所有加载和存储指令都表示为相对于该基地址的32位偏移量（最大4GB范围）。

以下是重写的具体示例：

  • 普通加载指令的重写：不安全的64位加载指令 ldr x0, [x1] 会被重写为安全的指令 ldr x0, [x27, w1, uxtw]。这里，x27 是保留的基址寄存器，w1 是32位偏移量。
  • 复杂指令的处理：对于像 ldp（加载寄存器对）这样不支持上述寻址模式的指令，需要引入额外的指令和第二个保留寄存器（如 x28）。首先通过掩码操作将地址加载到 x28，然后从 x28 加载数据，因为 x28 始终保证包含沙盒内的有效地址。

控制流方面，间接分支也采用类似的掩码操作来确保目标地址在沙盒内。Arm架构的定长指令编码是一个优势，因为无法跳转到指令中间，硬件会捕获此类错误。

对于系统调用等指令，会将其重写为从特定只读页加载函数入口点，然后分支到外部运行时代码来处理。线程本地存储也通过保留寄存器或运行时调用来虚拟化。

运行时与x86架构支持

上一节我们介绍了AArch64的实现，本节中我们来看看运行时环境以及x86架构的不同挑战。

LFI运行时是沙盒外部的可信代码，负责处理系统调用。它实现了Linux API的一个子集，这对于许多需要沙盒化的库（如解码器）来说已经足够，它们通常只需要分配内存和管理缓冲区。

将上述组件组合起来，使用流程如下：将软件编译为特殊目标，然后通过 lfi-run 工具在命令行运行，或者将沙盒库链接到应用程序中。

x86架构的支持面临不同的挑战：

  • 变长指令与捆绑：x86使用变长指令，可以跳转到指令中间。为了防止这一点，LFI采用“捆绑”策略，强制每条指令属于一个32字节的“束”，并在必要时填充NOP指令。在每条间接分支指令前，会清零地址的低5位，确保跳转总是对准束的起始位置。
  • 内存访问优化：x86没有Arm那种32位+64位的寻址模式。但可以利用 GS 段寄存器前缀进行优化，这比使用单独的指令来清零地址高32位性能更好。这被称为“段寄存器优化”。

性能评估与未来应用

上一节我们讨论了不同架构的实现，本节中我们来看看LFI的性能表现及其潜在应用场景。

我们在SPEC等基准测试上评估了LFI的性能，测试了多种配置：

  • 完全沙盒：标准的沙盒模式。
  • 仅存储沙盒：允许沙盒读取外部内存，但不能写入或破坏外部内存（适用于某些威胁模型）。
  • 仅跳转沙盒：不沙盒化加载和存储，可与MPK（x86）或PORE（Arm）等硬件内存隔离机制结合使用。

在AArch64上，完全沙盒模式的平均开销约为5%，仅存储模式约为1%，仅跳转模式约为0.5%。在x86上，完全沙盒模式性能相似，但仅存储和仅跳转模式的开销略高，这主要源于控制流隔离（捆绑机制）带来的开销。

针对Android的特定用例（如音频解码器），我们也进行了测试。例如，对Opus音频解码器和libvpx（VP8/VP9解码器，包含大量手写汇编）进行沙盒化，开销大约在5%左右。LFI在处理包含大量手写汇编的媒体解码器方面具有独特优势。

微基准测试显示，LFI的沙盒域切换开销极低，仅为数十个CPU周期，远低于进程间切换所需的数千个周期。

LFI还有一些正在进行的、有趣的应用探索：

  • 浏览器JIT引擎：修改JIT编译器，使其生成符合LFI安全子集的代码，并通过动态验证器确保安全性，从而大幅减少可信代码基。
  • 内核设备驱动隔离：在内核环境中，无法使用独立进程。LFI可以作为隔离第三方设备驱动程序、减少权限提升攻击面的潜在解决方案。

总结

本节课中我们一起学习了轻量级故障隔离技术。我们探讨了当前内存安全解决方案的局限性，介绍了LFI作为一种高效、兼容性强、安全的进程内沙盒化方案的愿景。我们深入了解了其在AArch64和x86架构上的技术实现细节，包括指令重写、控制流安全和运行时环境。最后，我们查看了LFI的性能评估数据，并展望了其在浏览器JIT和内核驱动隔离等领域的潜在应用。LFI旨在成为开发者工具箱中的一个强大新工具，以较低的代价为现有原生代码提供强大的隔离保障。

058：编写高质量测试

在本教程中，我们将学习如何为MLIR编写高质量、清晰且可维护的测试。我们将通过对比“好”与“坏”的测试示例，深入理解MLIR测试指南的核心原则，并掌握编写有效测试的具体技巧。

概述：为何测试指南至关重要

上一节我们介绍了本课程的目标。本节中，我们来看看测试为何如此重要。测试是我们代码库的一部分，它们为代码提供了文档。我们应该像对待产品代码一样认真编写和维护测试。

以下是测试的几个关键作用：

  • 文档功能：测试是理解特定转换或功能如何工作的绝佳资源。
  • 可发现性：我们需要确保测试易于查找和理解。
  • 一致性：一致的测试风格可以节省开发者的时间，避免他们在命名、格式等方面重新发明轮子，同时也让扫描测试以理解测试内容和边界条件变得非常容易。

核心原则：聚焦与简洁

现在，让我们通过一个具体例子来理解第一个核心原则。这个例子来自测试指南，展示了同一测试的“好”版本和“坏”版本。

输入MLIR代码（两个版本相同）：

func.func @test() -> index {
  %c1 = arith.constant 1 : index
  %c2 = arith.constant 1 : index
  return %c1 : index
}

这段代码定义了一个函数，其中常量1被定义了两次。测试会运行公共子表达式消除（CSE）来删除其中一个常量。

“坏”的检查指令：

// CHECK: func.func @test() -> index
// CHECK:   %[[C1:.*]] = arith.constant 1 : index
// CHECK:   return %[[C1]] : index
// CHECK: }

“好”的检查指令：

// CHECK: func.func @test()
// CHECK:   %[[C1:.*]] = arith.constant 1
// CHECK:   return %[[C1]]

这两个版本看起来相似，但“坏”版本存在几个问题：

  1. 检查了结尾的右花括号}，这对于测试转换是不必要的。
  2. 检查了索引类型index，但对此转换而言，类型并不相关。
  3. 在“好”版本中，我们使用了%[[C1:.*]]这样的变量捕获和块替换，使检查更清晰、更健壮。

这个例子虽然简单，但意义重大。当测试文件包含成千上万行时，保留不必要的内容会影响可读性和维护性。我们的测试应该只关注转换本身产生的影响。

避免重复与提供上下文

接下来，我们看看如何处理多个相似测试用例。这是一个来自真实测试文件的例子，测试将向量收缩（vector.contract）转换为外积（outer product）的转换。

最初，文件中有两个测试用例，它们看起来不同，但差异并不明显，难以一眼看出各自测试的边界条件。经过仔细比较，发现两个测试中定义访问模式的affine_map虽然变量名不同，但实际完全一样，这导致了不必要的重复。最终，其中一个测试被移除，因为它们在检查指令上没有体现出任何有意义的差异。

这个例子给我们的启示是：

  • 避免重复：重复的测试浪费资源且令人困惑。
  • 提供注释：如果你的测试用例有特殊之处（例如测试特定形状、数据类型或边界条件），请添加注释说明。否则，后来者只能猜测测试的意图。

利用命名传达语义

命名是让测试清晰易懂的强大工具。让我们看另一个测试向量掩码加载（vector.maskedload）降级的例子。

原始版本（存在问题）：

// 函数和变量名信息量低
func.func @test_vector_mask_load_f32(%arg0: memref<4xf32>, %arg1: vector<4xi1>) -> vector<4xf32> {
  %v1 = vector.maskedload %arg0[%c0], %arg1, %pass_thru : memref<4xf32>, vector<4xi1>, vector<4xf32> into vector<4xf32>
  return %v1 : vector<4xf32>
}

原始版本的问题在于：

  • 函数内的向量变量被命名为%v1，重复了类型信息。
  • 函数名@test_vector_mask_load_f32和参数名%arg0、%arg1没有传达任何语义。
  • 测试意图不清晰。

改进后的版本：

// 命名清晰传达了上下文
func.func @mask_all_true_f32(%base: memref<4xf32>, %pass_thru: vector<4xf32>) -> vector<4xf32> {
  %all_true_mask = vector.constant_mask [4] : vector<4xi1>
  %result = vector.maskedload %base[%c0], %all_true_mask, %pass_thru : memref<4xf32>, vector<4xi1>, vector<4xf32> into vector<4xf32>
  return %result : vector<4xf32>
}

通过改进命名：

  • 向量被命名为%all_true_mask，明确指出掩码全为真。
  • 参数命名为%base（内存基址）和%pass_thru（穿透值），说明了它们的角色。
  • 函数名@mask_all_true_f32直接点明了测试场景：当掩码全为真时，maskedload应退化为普通load。

清晰的命名让我们立刻理解了测试目的，并很容易发现测试覆盖的缺口。基于此，我们可以自然地补充更多边界测试用例：

// 补充测试：掩码全为假
func.func @mask_all_false_f32(%base: memref<4xf32>, %pass_thru: vector<4xf32>) -> vector<4xf32> {
  %all_false_mask = vector.constant_mask [4] : vector<4xi1>
  // 应直接返回 %pass_thru
  %result = vector.maskedload %base[%c0], %all_false_mask, %pass_thru : memref<4xf32>, vector<4xi1>, vector<4xf32> into vector<4xf32>
  return %result : vector<4xf32>
}

// 补充测试：混合掩码（负测试，模式不应触发）
func.func @negative_mixed_mask_f32(%base: memref<4xf32>, %pass_thru: vector<4xf32>) -> vector<4xf32> {
  %mixed_mask = vector.constant_mask [4] : vector<4xi1>
  // 此场景下转换不应发生
  %result = vector.maskedload %base[%c0], %mixed_mask, %pass_thru : memref<4xf32>, vector<4xi1>, vector<4xf32> into vector<4xf32>
  return %result : vector<4xf32>
}

在负测试的函数名中使用negative_前缀，可以清晰地表明这是一个期望转换不发生的测试。

工具使用与人工审查

编写测试，特别是生成检查指令（CHECK lines），可能是一项繁琐的工作。幸运的是，我们有自动化工具来辅助。

例如，可以使用Python脚本从MLIR输出自动生成初始的检查指令。这是一个很好的起点，可以节省大量手动输入的时间。

但是，工具生成的结果只是起点，而非终点。 开发者有责任对生成的检查指令进行审查和精炼：

  • 删除冗余：确保只检查转换必需的部分，移除无关的类型、符号或格式。
  • 添加上下文：使用有意义的变量名（如%[[RESULT:.*]]）和必要的注释。
  • 保持聚焦：检查指令应像测试代码一样清晰、简洁。

请务必使用自动化工具，但也请务必仔细审查它的输出。

总结与行动号召

本节课中，我们一起学习了MLIR测试指南的核心内容。我们了解到，编写高质量的测试需要做到聚焦简洁、避免重复、通过命名和注释提供清晰语义，并善用工具但不忘人工审查。

最后，我们发出一个行动号召：请帮助我们共同维护测试指南的质量。无论是您自己编写测试，还是在代码审查中看到他人的测试，都请确保它们符合本指南的要求。在审查时，请指出不遵循指南的测试，并请贡献者相应地更新它们。一致的、高质量的测试套件对整个社区都有益。

谢谢。

059：通过编译器锁孔窥探迁移至Clang

在本教程中，我们将学习如何将一个现有的代码库从GCC工具链迁移到Clang/LLVM工具链。我们将以Pixel Buds Pro第二代充电盒的迁移为例，探讨在迁移过程中遇到的关键挑战，特别是当目标平台（如RISC-V）包含自定义指令和重定位时，如何在不直接访问源代码的情况下，通过“锁孔”协作的方式解决问题。

概述：迁移的挑战与背景

上一节我们介绍了本次迁移的背景。本节中，我们来看看迁移面临的核心挑战。

我们之前已经成功将第一代Pixel Buds Pro从GCC迁移到了Clang。然而，第二代充电盒的迁移带来了更大的挑战，主要区别在于其核心从ARM切换到了RISC-V。这个RISC-V核心包含了供应商扩展，特别是自定义的分支指令，这些指令需要自定义的重定位。

不幸的是，供应商的工具链默认启用了这些扩展，且无法禁用。由于LLVM和LLD（LLVM链接器）目前不支持处理这些自定义重定位，我们无法将供应商的SDK与我们的代码链接起来。虽然LLD中有一个支持自定义重定位的PR正在审核，但尚未合并。

为什么这是个问题？在传统的操作系统（如Linux）中，应用程序和操作系统之间有明确的ABI（应用程序二进制接口）边界。然而，在微控制器（MCU）系统中，通常没有这样的边界。开发者需要将供应商提供的SDK（包含驱动、RTOS等）直接与应用程序代码静态链接，以优化二进制大小。这就要求应用程序和SDK之间的ABI必须兼容，而由于自定义重定位的存在，目前无法实现。

潜在解决方案与协作模式

上一节我们明确了问题所在，本节中我们来看看有哪些潜在的解决方案。

我们考虑了以下几种方案：

  1. 继续使用供应商工具链：这是我们最初的做法，但产品团队希望为耳机和充电盒使用统一的Clang工具链。
  2. 重写供应商SDK：重写所有使用自定义指令的代码以避免它们。但这没有现成的解决方案，且工程量巨大。
  3. 使用我们的Clang工具链重建供应商SDK：这需要供应商的支持。幸运的是，供应商非常配合，决定与我们合作迁移他们的代码库。但有一个关键限制：SDK代码是他们的专有IP，我们无法直接查看代码。

因此，整个协作模式就像是通过“锁孔”工作：

  • 我们构建一个Clang工具链（通常是上游最新版本，偶尔包含一些进行中的补丁或自定义Clang-Tidy检查），并将其发送给供应商。
  • 供应商使用该工具链构建新版本的SDK，链接应用程序，运行所有测试，然后将日志发送给我们。
  • 我们分析日志，尝试解决出现的问题，然后重复这个过程。

迁移策略：编译、优化与运行

上一节我们确定了协作方式，本节中我们来看看迁移代码库的一般策略。

迁移通常遵循三个步骤：首先让它能够编译，然后优化其大小，最后确保它能正确运行。在实践中，这些步骤往往是迭代和交叉进行的。

第一步：让它能够编译

即使Clang/LLVM努力支持大多数GCC扩展，它们也并非直接的替代品，仍然存在需要解决的重大差异。

以下是编译阶段常见的挑战：

  • 不同的警告：Clang可能会触发GCC未触发的警告。
  • LLVM libc的限制：在裸机（bare-metal）构建中，我们有意省略了某些接口（如pthread），因此有时需要修改代码以避免使用某些头文件。
  • 链接器（LLD）的诊断信息：LLD的错误信息有时不够友好。例如，错误信息“section can’t have both LMA and loaded row region”没有指出问题出现在链接脚本的哪一行、哪个段，这使得在大型链接脚本中定位问题非常困难。我们已提议改进LLD的链接脚本解析器。

第二步：让它能够适配（优化大小）

编译成功后，下一步是优化二进制大小，这对资源受限的MCU至关重要。

许多MCU使用异构内存布局，这意味着内存不是统一的，符号的放置位置对正确性和性能都很重要。然而，标准的C/C++后段（backend）并不直接支持异构布局，开发者通常需要借助非标准扩展。

以下是一个处理异构内存布局的常见模式及其问题：
假设我们想将代码放入一个名为retain_ram的特定内存区域。我们有几个不同的符号（变量和函数）。在裸机系统中，通常会使用-ffunction-sections和-fdata-sections编译，以便链接器可以进行垃圾回收（GC），移除未使用的代码。每个符号会进入自己的段（section）。

为了便于管理，开发者通常使用section属性来注解所有应放入retain_ram的符号，然后在链接脚本中手动放置整个.retain_ram段。但这样做有一个大缺点：它将所有符号集中在一个段里，破坏了链接器的垃圾回收功能，因为链接器无法丢弃该段中未使用的单个符号。

我们最初的解决方案是编写一个Clang-Tidy检查，自动将每个使用section属性的符号拆分到独立的、名称唯一的段中（例如.retain_ram.foo， .retain_ram.bar）。但这非常侵入式，供应商代码库中有数千个section属性实例，难以维护。

更好的解决方案：我们最终在LLVM中推广并实现了一个新特性：-fseparate-named-sections。启用此标志后，每个被section属性注解的符号都会获得自己独立的段（但段名相同）。这恢复了链接器的垃圾回收能力，且无需修改源代码。其原理是，在ELF格式中，允许存在多个同名的段。

不过，这种方法无法完全恢复所有语义。例如，对于应进行零初始化（.bss）的变量，如果手动将其放入.retain_ram段，它们就会占用二进制空间。理想的解决方案可能是引入一种新的语法来更好地表达“将此符号放入特定内存区域”的语义，并配合新的链接脚本语法。这仍在讨论和开发中。

为了进一步缩小体积，我们还尝试了链接时优化（LTO）。我们最初尝试用-flto=thin编译运行时库，但发现LTO与运行时库的配合存在一些问题。目前，我们仅对应用程序代码使用LTO，并希望未来能扩展到整个应用。

第三步：让它能够运行

即使代码编译成功且大小合适，也可能无法正确运行。我们遇到的主要问题是性能。

该应用有严格的时序要求（涉及蓝牙通信）。性能问题的根源在于内存访问对齐。虽然该MCU支持非对齐读写，但基准测试显示其速度比对齐读写慢大约4倍。由于原始代码依赖一些未定义行为，而GCC和Clang在处理这些未定义行为时存在差异，导致Clang生成的代码将某些符号放在了非对齐地址上，从而引起性能下降。

我们使用未定义行为消毒剂（UBSan） 来定位所有非对齐访问。我们使用了-fsanitize=alignment，并搭配一个比现有版本更精简、且与我们的RISC-V架构兼容的自定义运行时库。未来我们希望能在上游直接支持这种轻量级的嵌入式UBSan。

我们还遇到了其他性能瓶颈。开发者并不总是能准确预测代码性能，例如可能将热函数（hot function）放在源文件（.c）中，导致编译器无法跨编译单元内联；或者将冷函数放在头文件中标记为inline，不必要地增加了代码体积。

为了解决这些问题，我们采用了基于插桩的性能分析（PGO）。我们使用标准的插桩标志编译代码，通过JTAG连接设备运行程序并收集性能分析数据，然后使用优化备注（optimization remarks）为开发者提供改进代码的指导。例如，编译器可能会提示“无法内联函数foo，因为看不到其定义（但它很热，应该被内联）”，开发者据此将函数移到头文件中并标记为static inline。未来我们希望自动化这个过程。

最后，在调试此类问题时，调试器非常有用。我们成功使用了LLDB。但当应用崩溃时，我们常常会陷入像中断处理程序这样的地方，而这些程序通常用汇编编写。为了获得良好的堆栈回溯，需要编写正确的DWARF调用帧信息（CFI）注解，但这非常容易出错。

这促使我们提议并开发了一个DWARF CFI验证器原型。它是一个简单的抽象解释器，可以运行在你的CFI注解上，尝试发现并标记问题。目前原型仅支持x86_64，但我们正计划将其扩展到ARM或RISC-V等嵌入式领域更常见的架构。

总结与社区邀请

本节课中我们一起学习了将一个嵌入式代码库从GCC迁移到Clang/LLVM的完整过程，特别是在无法直接访问供应商SDK源代码的“锁孔”模式下所面临的独特挑战和解决方案。

回顾整个努力，有人可能会问：为单个应用程序做这一切是否值得？对我们而言，答案是肯定的。原因如下：

  1. 在嵌入式领域，硬件平台的生命周期较长，同一款MCU可能用于多代产品。投入精力使其与最新的编译器工具链协同工作是值得的。
  2. 在此过程中，我们提出了许多新特性和改进的想法，其中一些已经融入LLVM，另一些有望在未来实现。
  3. 我们开发的所有自动化工具不仅适用于本项目，也对其他嵌入式项目乃至整个社区有益。

如果你对嵌入式开发感兴趣，欢迎加入LLVM内部日益壮大的嵌入式社区。我们有一个公开的工作组，每四周举行一次会议（太平洋时间周四上午9点）。我们还会举办研讨会。期待你的参与！

问答环节

问： 你提到供应商SDK迁移到了你们的LLVM工具链，这包括所有工具吗？还是只包括编译器？
答： 包括所有工具。我们使用与所有项目相同的工具链，即直接从LLVM Git仓库构建的Vanilla LLVM。我们使用Clang作为编译器，使用所有LLVM工具，使用LLD作为链接器，并且使用LLVM libc作为C库，libc++作为C++库。我们与LLVM libc团队（包括Michael）密切合作，他们非常出色地解决了我们遇到的各种问题，例如缺失的函数、内存拷贝（memcpy）的性能改进等。这次协作至关重要。

问： 这个设备里大概有多少行代码？
答： 实际上我不知道具体的代码行数，因为我们通常不以代码行数来衡量。但我知道二进制大小，这才是我们真正关心的。对于Pixel Buds Pro第二代充电盒，镜像文件大约在1MB左右。

问： 与供应商的互动有多困难？说服他们尝试新编译器容易吗？
答： 互动非常积极。我们很幸运，在供应商那边找到了一位非常乐于合作的工程师。最具挑战性的方面是沟通延迟，因为他们不在美国，位于不同的时区。此外，他们将新工具链集成到自己的工作流程中也需要时间。因此，整个项目持续了一年多，但大部分时间是等待反馈，并非全是主动开发时间。

问： 关于你添加的链接器标志（-fseparate-named-sections），它是否完全尊重指定的段名？你是在ELF中放置多个同名的段，还是像-fdata-sections那样附加变量名？
答： 使用-fseparate-named-sections时，我们完全不修改段属性。我们完全尊重属性中指定的内容，只是在ELF中生成多个同名段。

问： 我的客户有很多汇编代码，他们抱怨调试困难，因为他们不写DWARF信息。你提到的DWARF CFI验证器何时能在上游可用？
答： 我展示的功能已经在上游了。目前有一个开放的PR（尚未合并）改进了我们能检测到的问题类型。最大的挑战是，这项工作是由我们团队去年夏天的一位实习生完成的，实习期结束后需要找到其他人愿意接手继续开发。另一个问题是，我们在MC层缺少足够的语义信息来真正理解单个指令的作用（例如对堆栈的影响）。我们发现这些信息已经存在于LLVM的其他地方（例如LLVM的MCPlus），但将其引入到LLVM中以供验证器使用将是一项重大的重构工作。

问： 我们尝试用Clang编译GCC编译的代码时，遇到过一些Clang无法编译的结构，后来发现这些根本不是有效的C++代码，只是GCC有自己的（可能是错误的）解释。你们遇到过类似情况吗？
答： 我们偶尔会遇到类似问题。一个非常具体的、几乎在每个嵌入式代码库中都能看到的例子是optimize属性。这个属性很流行，它允许你告诉编译器仅针对特定函数改变优化级别。Clang/LLVM不支持这个属性，而且基于LLVM优化管道的工作方式，我认为我们永远不会支持。解决方法是改变代码组织方式，例如在编译单元（P）级别进行设置。我认为整理一份“常见迁移问题与解决方案”手册会非常有益，或许可以发布在LLVM网站上，同时提供更多像Clang-Tidy这样的自动化工具来帮助开发者解决这些问题。

060：性能优化实践与未来方向

概述

在本节课中，我们将学习如何为爱立信（Eson）的自研多核架构（MC）构建MLIR编译器，以实现实时AI推理。我们将重点探讨在追求低延迟和低功耗的约束下，如何通过编译器优化提升性能，以及在此过程中遇到的挑战和获得的经验教训。

上游优先的开发策略

我们采用了“上游优先”的开发策略，即优先将代码贡献到LLVM/MLIR等开源项目的主线。这一策略带来了显著的开发效率提升。

以下是该策略带来的主要优势：

  • 快速原型验证：重用从Linalg到LLVMIR的现有方言，使我们能在约两周内让模型在实际硬件上运行。
  • 无缝工具链集成：与现有的基于LLVM的工具链实现了良好的集成。
  • 免费优化：能够直接利用上游社区开发的多种优化技术。

卷积优化的尝试与取舍

上一节我们介绍了上游优先策略的优势，本节中我们来看看具体的优化尝试。我们尝试了不同的卷积优化方法。

以下是我们的具体发现：

  • Im2col优化：这种流行的优化方法对我们的硬件并不有利。因为我们并非针对专门的AI加速器，而是普通核心，无法从中受益于其内存访问模式。
  • Winograd优化：这种方法为我们带来了不错的性能提升。然而，Winograd算法在数值上并不稳定，我们遇到了模型崩溃的问题。在后续尝试量化时，问题变得更加严重。因此，我们不得不暂时搁置这项技术。

探索指令级并行（ILP）

对于我们的硬件而言，性能提升的“北极星”是指令级并行（ILP）。向量化是实现ILP的一种自然思路。

我们选择了仿射向量化（Affine Vectorization），而非循环向量化（Loop Vectorization），因为后者对一些我们想运行的模型支持不足。但遗憾的是，启用向量化后，我们并未看到性能提升，反而观察到了三倍的性能下降。

经过深入分析，我们发现了两个关键问题：

  1. 向量化维度选择不当：仿射向量化器总是选择最内层循环进行向量化。对于卷积而言，这通常是卷积核维度，其循环次数（trip count）很低。这导致生成的向量操作中，大部分通道（lane）都被掩码（mask）关闭，效率低下。理想的向量化维度应该是图像维度。
  2. 后端支持缺失：我们的后端并未对掩码向量操作（masked vector ops）进行优化处理。这些操作被完全展开成了大量的条件判断语句，导致单个向量操作可能膨胀为上百条指令。

这个教训非常重要：优化并非免费的午餐。一项优化本身不会自动带来性能提升，必须确保它以合理的方式应用，并且具备必要的底层支持机制才能发挥效果。

融合：更有效的ILP驱动因素

尽管向量化的尝试受挫，但这推动我们寻找其他提升ILP的方法。我们发现，对于我们的硬件，操作融合（Fusion）是比向量化更有效的ILP驱动策略。

我们可以与模型开发者协作，共同创造更好的ILP机会。例如，将标准卷积替换为可分离卷积（separable convolution）。这使得线性元素级操作融合能够更有效地工作，因为它可以垂直方向进行融合，而不再受限于归约迭代器。

通过这种方式，我们在其他方面等效的模型上，观察到了超过100倍的性能提升。

量化的必要性与协同设计

对于我们的DSP核心，它们为定点计算而非浮点计算进行了深度优化。因此，量化对我们至关重要。

然而，最好的量化策略是从模型侧开始构建。在DSP聚焦的系统中，输入数据很可能已经是量化的。模型开发者可以利用量化感知训练（QAT）技术，在最大化精度和性能的同时进行量化。

通过将这些经验教训传达给模型开发者，我们实际上在很大程度上绕开了编译器层面繁重的性能优化工作，通过软硬件协同设计取得了显著成效。

未来方向

最后，我们展望一下未来的工作方向。

以下是几个我们正在探索的重点领域：

  • 超低比特量化：行业趋势正朝向4比特或更低的超低精度量化。这类操作在我们的硬件上原生支持不佳，解包操作ILP极低，会导致严重的性能回退。我们正在研究如何启用自动SIMD，以不解包的方式直接进行计算。
  • 加速器集成：探索如何让模型开发者通过AI框架（如JAX）有效地利用我们的DSP和定制加速器。
  • 全栈可观测性：在整个MLIR栈中启用可观测性，包括IR快照、诊断信息、变换过程的可追溯性，并确保其与现有的追踪基础设施有效集成。我们希望很快能在上游社区完善这些功能。

总结

本节课中我们一起学习了为爱立信多核架构构建MLIR编译器的实践经验。我们认识到上游优先策略能加速开发，但优化技术的选择必须结合硬件特性。我们经历了向量化带来的性能陷阱，发现了操作融合是提升ILP的更有效途径，并强调了模型与编译器协同设计（尤其是量化）的重要性。最后，我们展望了应对超低精度量化和提升系统可观测性等未来挑战的方向。

061：使用libc进行链接时间优化的历险记

概述

在本节课中，我们将要学习在LLVM工具链中，将标准C库（libc）作为源码参与链接时优化（LTO）所遇到的挑战和解决方案。我们将探讨LTO的基本流程、遇到的问题、以及如何通过更精细地控制编译器的“内置函数”处理机制来安全地实现优化。

LTO流程简介

上一节我们介绍了课程主题，本节中我们来看看链接时优化的标准流程。大多数关注代码体积的项目更倾向于使用完全LTO，因为它在体积优化上通常比瘦身LTO效果更好。

完整的LTO流程如下：

  1. 链接器接收一组LLVM IR文件，进行符号解析。
  2. 控制权移交到LTO管线，开始合并模块并进行内部化。
  3. 最终得到一个近乎完整的、单块的LLVM IR程序。
  4. 对该IR进行优化和代码生成。
  5. 常规链接过程从此点恢复，最终生成输出文件。

我们的核心动机是最大化编译器的优化能力。LTO通过合并模块、模糊由独立编译单元引入的边界，为编译器优化接口、删除无用代码提供了绝佳机会。

遇到的问题：符号消失之谜

上一节我们介绍了LTO的理想目标，本节中我们来看看实践中遇到的第一个问题。如果你尝试对libc源码使用LTO，传递 -flto 标志，编译过程可能会失败。

你可能会遇到一个关于 bcopy 的未定义隐藏符号错误。这很奇怪，因为 bcopy 明明是你编译的源码的一部分，它就在 libc.a 库里。

为了探究原因，我们可以使用 -save-temps 标志来查看优化管道中发生了什么。你可能会注意到，bcopy 在IR中确实有定义。那么，为什么链接时会说它未定义呢？

让我们通过一个简化示例来深入理解。假设我们有一个函数 foo，它包含一个条件分支，分支中调用了 memcpy。

define void @foo(...) {
  ...
  call void @memcpy(...) ; 外部符号
  br i1 %cond, label %true, label %false
  ...
}

同时，模块中也有 bcopy 的定义。在全局死代码消除（Global DCE）运行后，bcopy 没有被引用，因此被编译器删除。这看起来是合理的优化。

然而，随后 simplify-libcalls 优化会运行。这个优化可能会将 memcpy 调用转换为 bcopy 调用，因为它在分支条件中被使用。但此时，bcopy 在模块中只剩下一个外部引用声明，其定义已在之前的DCE阶段被删除。因此，当链接最终完成时，就会因为找不到 bcopy 的定义而失败。

总结问题原因：

  1. libc的所有API在LTO开始时被标记为内部（internal）。
  2. 全局DCE发现 bcopy 未被使用，将其删除。
  3. simplify-libcalls 优化将 memcpy 调用转换为 bcopy 调用。
  4. 链接失败，因为 bcopy 已不存在。

根本原因：编译器对libc的假设

上一节我们看到了一个具体的失败案例，本节中我们来分析其背后的根本原因。使用 libc 作为源码与常规使用方式有本质区别。

传统编译器对libc有一系列假设，这些假设基于静态或动态链接模型。编译器将libc接口视为一个抽象的、保证提供的接口，从而可以进行高级转换，例如将 memcpy 模式匹配并优化为 bcopy。

控制这种行为通常通过 -fno-builtins 等标志实现。此外，LLVM后端也有一套用于处理库调用和内置函数的低级转换，这些转换无法单独关闭。

初步解决方案：使用运行时库调用

上一节我们分析了问题的根源，本节中我们来看看第一个想到的解决方案。运行时库调用机制看起来很有吸引力，例如 memcpy 就是这样处理的。

这些调用在优化管道中被区别对待：

  • 它们不能被删除，因为预期可能由外部提供定义。
  • 它们总是从比特码归档文件中提取，因此始终可用。

我们最初的尝试是直接使用这个机制。但在实施前，我们了解到这并非最佳方案。我们需要在概念上区分两个阶段：

  1. 允许使用内置函数的世界：例如 memcpy、bcopy。此时libc被视为完全抽象的接口，LLVM通过 -fno-builtins 等机制使用它。
  2. 无内置函数的世界：进入此阶段后，优化器将整个程序视为一个整体，可以自由进行任何它认为合适的变换。

我们需要在IR中编码这种区分，并在管道中找到一个安全点来切换。

实现更精细的控制

上一节我们讨论了概念模型，本节中我们来看看如何具体实现。一个天真的想法是：在管道的某个安全点，简单地为模块中的所有函数添加 nobuiltin 属性。

我们构建了这样一个Pass。它确实有效，因为 bcopy 不再被DCE删除，memcpy 也不会被重写为 bcopy。

然而，这种方法副作用很大。无论你在管道的哪个位置添加 nobuiltin 属性，它都会显著改变编译器的行为，影响后端通过目标库信息（TLI）进行的许多优化。例如，原本可以将 sqrtf 调用替换为单条X86指令的优化将无法进行。这相当于“把婴儿和洗澡水一起倒掉了”。

构建安全的后期函数使用规则

上一节我们看到了简单方案的局限性，本节中我们将介绍一套更安全、更精细的规则。我们意识到，问题在于链接完成后，我们需要对如何处理这些符号有更精确的认识。

我们制定了一套规则，用于安全地在链接完成后使用函数：

以下是安全使用函数的核心规则：

  1. 允许添加对已链接比特码的引用：如果一个符号（如 bcopy）因为某种原因被拉进了链接，那么将 memcpy 转换为对它的调用是安全的。我们需要允许这类安全的转换发生。
  2. 保护已链接的函数：为了安全地进行上述转换，被转换的目标函数必须被视为“神圣不可侵犯”。LLVM已有机制实现这一点：将其标记为外部（external）链接，并防止其被内部化。这样，优化通道就不能删除它或更改其ABI。
  3. 禁止添加对新函数的引用：链接完成后，哪些libc函数在链接中、哪些不在，这个集合就冻结了。不允许在编译过程中途再决定引入新的比特码并重新进行内联等操作。如果链接时 bcopy 没被需要，那么之后即使发现可以将 memcpy 优化为 bcopy，也不允许这样做。我们认为这是一个合理的取舍。

基于这些规则，我们构建了一个补丁。令人惊讶的是，实现起来并不复杂，只有几百行代码，且对现有代码的修改面广但改动点小。

解决方案的工作流程

上一节我们介绍了安全规则，本节中我们详细看看解决方案的具体工作流程。

以下是该方案的关键步骤：

  1. 链接器收集信息：LLD需要知道哪些是libc函数。我们可以借鉴已有的 getRuntimeLibcall 函数机制。通过目标三元组查询目标库信息（TLI），可以获得一个保守的libc函数列表（如 memcpy， bcopy， fmemcpy...）。
  2. 链接器解析符号：链接器遍历这个列表，查看这些符号解析到了什么。确定哪些是以比特码形式存在的（这些是潜在的问题点），哪些是手写汇编实现（无需特殊处理）。
  3. LTO侧确定可用集合：在LTO阶段，扫描接收到的所有IR（完全LTO）或模块摘要（瘦身LTO），确定哪些libc函数通过直接调用等方式实际进入了项目。这个集合将被冻结。
  4. 构建受限的TLI：在LTO过程中调用代码生成和优化时，负责构建TLI对象。此时，可以根据上一步的“可用集合”，禁用那些未进入链接的函数的转换，并将已进入链接的函数标记为外部链接以保护它们。

这个方案的一个重要优点是：除非你确实在尝试对libc进行LTO，否则它完全是一个空操作，这增加了其被上游接受的可能性。

未来方向与社区协作

上一节我们介绍了一个可行的技术方案，本节中我们展望未来的工作并呼吁社区协作。我们还有一些评估工作要做，并希望能在Google内部更广泛地测试。

我们正在几个嵌入式项目中进行调查，目前效果良好，获得了不错的代码体积改进。我们相信通过调整还可以做得更好。

未来需要与上游社区进行大量讨论：

  • 是否需要为libc函数添加属性以使方案更完善？
  • 是否应该修改Pass管道？
  • 是否应该以更原则性的方式在IR中编码这些信息？
  • 我们是否也想对内置函数（builtins）进行类似处理？
  • 我们需要明确LTO在这些情况下的确切语义，以及它与传统链接模型的区别。

我们无法独自完成这一切，期待社区成员的思考和参与。我们将在相关讨论中发布更新和补丁链接。

总结

本节课中我们一起学习了在LLVM工具链中对libc源码使用链接时优化所面临的挑战。我们从LTO流程和遇到的问题出发，分析了编译器对libc的传统假设为何在源码LTO场景下失效。接着，我们探讨了从简单的运行时库调用机制到更精细的“无内置函数”世界切换方案，并指出了其局限性。最后，我们介绍了一套安全的后期函数使用规则，以及基于此实现的一个具体工作流程，该流程通过在链接后冻结libc函数集合并保护已链接函数，使得对libc的LTO成为可能。我们认识到这只是一个开始，未来需要在性能权衡、机制设计和社区协作上继续努力，以充分释放编译器在源码LTO场景下的优化潜力。

062：探索性能优化机会

概述

在本节课程中，我们将探讨如何通过专业化MLIR的数据结构来提升编译性能。我们将分析两种主要的权衡：可扩展性和可用性，并评估通过限制某些MLIR特性所能获得的潜在性能收益。核心思想是：如果MLIR的数据结构是一个接口，允许我们为特定编译器替换更优的实现，我们能获得怎样的性能提升？

可扩展性的权衡

MLIR的一个主要优势是其强大的可扩展性，能够表示机器学习、硬件设计、计算机图形等多个领域的中间表示。然而，这种通用性带来了运行时开销，进而影响了编译时间。

操作的内存布局需要更大以支持可变数量的操作数和结果。
检查操作是否实现了某个接口或特质需要间接层，存在开销。
操作码检查和属性字典的使用也带来了额外成本。

可用性的权衡

与LLVM类似，MLIR提供了高效的IR操作算法，例如快速替换、插入和移动操作。这些算法的实现依赖于指针链接，维护这些链接的正确性需要大量的簿记工作。

例如，替换一个操作的汇编代码量相当可观，而直觉上这可能只是简单地设置一次内存。

核心问题：数据结构作为接口

那么，我们能否在保留MLIR可扩展性的同时，获得类似LLVM这种具有固定IR的即时编译器的性能？我们能否同时拥有高效的渐近算法，并在愿意以不同方式编写编译器通道的前提下，使某些通道运行得极快？

本质上，我们提出的问题是：如果MLIR的数据结构是一个接口，允许我们为当前编译器替换更优的定义，会怎样？

性能优化机会

以下我们将通过几个微基准测试，探索专业化数据结构可能带来的性能提升。

操作码检查优化

目前，MLIR使用dyn_cast/cast机制来检查操作类型。我们可以实现一个类似的机制，为每个操作分配一个固定的操作码，从而减少内存使用并加速检查。

性能对比：

  • 迭代操作向量：0.31纳秒/操作
  • MLIR当前检查（如检查是否为add操作）：0.51纳秒/操作
  • 使用固定操作码检查：0.29纳秒/操作

在缓存命中的理想情况下，有显著提升。在非缓存场景下，仍有约10%的性能收益。

接口访问加速

目前访问接口需要遍历操作的所有特质和接口，速度较慢。如果改用基于switch-case的实现，理论上可以获得4倍的加速，达到LLVM在此类操作上的性能水平。

减少内存占用

通过分析操作的内联结构，我们可以移除编译器特定场景下不需要的特性，以减少内存占用。

可移除或缩小的部分：

  • 属性字典：如果编译器已知所有操作属性，可使用操作属性替代。
  • 操作名称：如果操作数量有限，可将其编码为16位或8位整数。
  • 支配索引：如果通道不使用支配关系或验证，可移除。
  • 位置信息：发布版本中可移除调试位置信息。

通过上述优化，操作创建速度可获得10%-20% 的提升。在一个将常量置于右侧的规范化通道中，处理十亿级操作时，时间可从7.1秒降至6.1秒。

固定操作属性大小与分配重用

在MLIR中，操作属性大小差异很大。如果我们已知所用方言的操作属性大小固定且很小，可以将其内联分配到每个操作中。

这带来了一个关键优化：操作分配重用。

传统操作替换流程：

  1. 为新操作分配内存。
  2. 初始化新操作。
  3. 替换所有使用旧操作结果的地方（检查每个用户，成本高）。
  4. 删除旧操作。
  5. 释放旧操作内存。

分配重用优化流程（当操作数列表相同时）：

  1. 重用旧操作的内存分配。
  2. 更改操作码字段。
  3. 删除旧操作数。
  4. 就地初始化新操作（例如设置新常量）。

在一个简单的常量折叠基准测试中：

  • 迭代IR：3.6纳秒/操作
  • 传统常量折叠替换：114纳秒/操作
  • 分配重用替换：~7.2纳秒/操作 （约为迭代成本的两倍）

此项优化带来了约15倍的加速。这只有在已知操作属性大小固定时才可能实现。

限制MLIR特性的权衡

上一节我们探讨了在可扩展性方面进行专业化的机会。本节中，我们来看看如果主动限制MLIR的某些特性，例如定值-使用链，能带来哪些性能提升。这类似于WebKit等超快编译器采用的策略。

移除定值-使用链

定值-使用链用于从操作结果获取所有使用该结果的地方，它通过链表实现。移除它可以减少内存访问和指针维护开销。

性能影响：
在创建具有多个操作数的操作时，移除定值-使用链可以显著降低开销。例如，创建具有七个操作数的操作，时间可以从77纳秒减少到45纳秒，提升明显。

在之前的“常量置右”规范化通道中，移除定值-使用链能带来额外的1.5倍加速。

需要注意的是，移除定值-使用链后，替换单个操作将变得复杂，可能需要同时替换一批操作或使用更复杂的算法来跟踪用户。

使用向量替代链表

另一个想法是将操作存储在连续向量中，而非链表中。这能进一步提升局部性和访问速度。

然而，这带来了巨大的代价：你无法在IR中任意位置高效地插入或替换单个操作，可能需要重写整个IR片段。这改变了MLIR的核心操作模型。

竞技场分配

当前，每个操作都通过malloc单独分配内存，每次分配都有成本，且每个指针占用64位空间。

我们可以采用竞技场分配策略：预分配一个大数组（竞技场），操作在其中分配，并使用数组索引而非直接指针来引用操作。

在假设所有操作大小相同的实验中：

  • 对于最多包含7个操作数的操作，创建IR的速度可提升4-5倍，重写IR（常量折叠）的速度提升约50%。
  • 对于最多包含3个操作数的操作，创建IR的速度可提升5倍，重写IR的速度提升2倍。

这种数组分配机制若能集成到MLIR中，可能带来显著的益处。

总结

本节课我们一起探索了通过专业化MLIR数据结构来优化编译性能的实验和思路。

主要发现：

  1. 内存优化收益显著：通过减少内存占用和固定数据结构布局，能在多个微基准测试中获得10%-20%的性能提升。
  2. 分配重用潜力巨大：在已知操作布局固定的前提下，重用内存分配可使特定操作替换获得数量级（15倍）的加速。
  3. 限制特性以换取速度：移除如定值-使用链等特性，可以进一步减少开销，但会改变编程模型，需要更复杂的算法支持。
  4. 竞技场分配有前景：采用基于数组的分配策略，能大幅提升操作创建和重写的速度。

这些实验表明，如果能够约束MLIR的使用范围或引入可配置的数据结构策略，存在强大的性能优化机会。未来，也许我们能在MLIR中提供一个LLVM方言，使其内部的LLVM通道运行得和原生LLVM一样快。同时，也可以考虑将其中一些优化（如可控的分配重用或竞技场分配）作为可选项暴露给MLIR用户，而不仅仅是停留在理论探讨层面。

063：为LLVM libc添加宽字符支持 🧩

概述

在本教程中，我们将学习一个名为“Widen Your Char-izons”的暑期项目。该项目旨在为LLVM的libc库添加宽字符转换支持。我们将从理解字符串在内存中的表示方式开始，逐步深入到宽字符与多字节字符的转换原理、接口设计以及具体的实现细节。本教程内容简单直白，适合初学者理解宽字符处理的基本概念。

字符串在内存中的表示

在C语言中，字符串通常表示为一个字符数组，其中每个字符是一个一字节的数据类型。这种表示方式对于单字节字符非常方便。

然而，当处理需要超过一个字节的字符时，例如表情符号或其他字母系统，这种表示方式就变得复杂。下图展示了字符串“Llvm”在内存中的表示，其中每个字符恰好占用一个字节。

多字节字符与宽字符

为了解决上述问题，引入了多字节字符和宽字符。它们都是用于表示超出ASCII范围的字符的方法。

多字节字符串通常采用UTF-8编码，其字符大小可在1到4个字节之间变化。因此，字符串的字节长度不一定等于其字符个数。它通过char指针引用，解析时有可能停在某个字符的中间。

宽字符字符串则通常采用UTF-32编码。它总是占用固定数量的字节（例如4字节），这使得计算字符串长度变得容易，并且不会在字符中间停止解析。

以下是字符串“clown σ”在多字节（UTF-8）和宽字符（UTF-32）下的表示示例。

字符转换过程

屏幕上展示了一个码点（UTF-32字符表示）与其等效UTF-8形式之间的转换表。从这个表可以得出的主要结论是：观察第一个字节的前几位，可以获知该字符在UTF-8编码中将占用多长。

利用这个信息，让我们通过一个具体例子，学习如何将多字节字符串转换为宽字符串。我们将尝试把字符串“aσ”从多字节转换为宽字符。

旁边这个“部分状态”将表示我们在转换过程中的进度。

  1. 首先，我们从输入字符串读取一个字节并放入部分状态。字母“A”的第一个字节表明它只占一个字节。此时，部分状态包含一个完整字符，我们可以将其写入输出宽字符字符串。
  2. 接着，我们读取下一个字符。其第一个字节表明该字符将占用两个字节。我们继续从输入字符串读取。
  3. 再次读取后，部分状态现在包含一个完整字符，我们将其写入输出宽字符字符串。至此，转换过程完成。

接口使用示例

现在，我们来看一个使用该接口的例子。函数mbrtowc用于将多字节字符转换为宽字符。

我们传入一个用多字节字符串表示的小丑表情符号作为输入，输出变量是宽字符字符串。初始调用时，我们指定读取1个字节。但由于小丑表情符号占4个字节，函数返回-2。

接着，我们传入读取3个字节（因为之前已读取1字节，总计需4字节）。此时，它成功读取了小丑表情符号，并返回3（成功读取的字节数），同时将宽字符放入输出字符串。

你可能注意到该函数还接受一个mbstate_t参数。这是因为mbrtowc是一个可重启版本的函数，用户可以通过mbstate_t参数获得更多控制权，并能在任何时候重启转换。不可重启的函数则维护一个全局的内部状态，可以从中断处继续转换，但无法强制其重启。

设计思路：mbstate_t 与字符转换器

我们反复提到了mbstate_t，它代表了转换过程中的部分状态。这是一个由libc本身定义的不透明结构体。作为LLVM libc的开发者，我们可以决定在这个结构体中包含什么内容。

我们最终决定存储以下内容：

  • 以UTF-32形式存储的部分字符。
  • 已存储到该部分状态的字节数。
  • 该字符预期占用的总字节数。

这使得检查部分状态是否包含完整字符变得非常简单，只需检查“已存储字节数”是否等于“预期总字节数”。

基于此，我们设计了字符转换器类，这是我们用来与mbstate_t交互的内部接口。它包含push和pop函数，以及其他如clear、is_empty、is_full和is_valid_state等功能。

  • 当从多字节转换为宽字符时，我们一次push一个字节，然后pop出整个宽字符。大部分转换逻辑发生在push过程中。
  • 当从宽字符转换为多字节时，我们push整个宽字符，然后一次pop一个字节。大部分转换逻辑发生在pop过程中。

设计思路：字符串转换器

在成功实现不同字符编码间的转换后，我们希望更进一步，实现不同字符串之间的转换。

为此，我们设计了字符串转换器类。它接收某种编码类型的输入源字符串，其接口允许我们持续地pop出字符（单字节或整个宽字符），直到整个源字符串转换完毕。这使字符串转换的整个过程变得更加容易。

最后，我们实现了内部的可重启函数，它们等同于参数最多的公共可重启函数。我们的两个内部字符转换函数被字符转换函数调用，而两个内部字符串转换函数则被所有三个字符串转换函数以及mbrlen（用于计算多字节字符串长度）调用。

关键设计决策

1. mbstate_t 的大小与存储内容

我们的最终决定是使用6字节：

  • 4字节用于保存部分转换结果。
  • 各1字节分别用于存储“预期总字节数”和“已存储字节数”。

我们考虑过其他方案：

  • 第一备选方案：仅用4字节保存部分转换。但这会带来问题，因为每次需要获取状态时，都必须推导出总字节数和转换状态。
  • 第二备选方案：利用宽字符只使用21位的事实，用剩余的位来存储总字节数、已处理位数和一些未使用的位。但此方案的问题是，如果宽字符未来扩展到超过21位，整个系统就需要修改。

2. 字符串转换器类的设计

最初我们并未计划设计这个类，只打算创建一个字符转换器类来处理所有转换。但随着时间推移，我们发现使用字符转换器类来实现公共的字符串转换函数，会使实现变得复杂且难以推理。

同时，我们也考虑到未来的可扩展性，例如最终需要添加UTF-16支持。因此，拥有一个易于使用的接口来转换不同字符串格式非常重要。最终，我们设计了字符串转换器类来提供这种灵活性。

未来工作展望

我们成功为LLVM libc添加了宽字符支持，但仍有很大的发展空间。

  1. 文件操作支持：目前尚未在文件操作中支持宽字符。
  2. 宽字符输出：实际上还不具备打印宽字符的能力。
  3. 平台差异支持：本项目专注于32位宽字符，但某些平台（如Windows）期望宽字符仅为16位。我们需要增加对此的支持。
  4. 实现 wctype.h：这是宽字符版本的ctype.h头文件，包含如iswalpha、iswdigit、iswupper等函数。
  5. 浮点数转换：我们实现了宽字符到整数的转换，但浮点数转换逻辑更为复杂，未能完成，而这将非常有用。

总结

本节课中，我们一起学习了“Widen Your Char-izons”项目，该项目为LLVM libc库添加了宽字符支持。我们从字符串的内存表示基础讲起，理解了多字节字符与宽字符的区别及其编码方式。通过实例分析了字符转换的过程，并探讨了可重启函数接口的使用。在核心设计部分，我们深入了解了mbstate_t状态的设计、字符转换器和字符串转换器类的职责与协作，以及项目中的关键设计决策。最后，我们展望了未来可能的改进方向。希望本教程能帮助你理解宽字符处理的基本概念及其在系统库中的实现思路。

064：如何在嵌入式应用中测试和评估 LLVM-libc 🧪

概述

在本节课中，我们将学习如何测试和评估 LLVM-libc，以判断其是否适合作为嵌入式应用的工具链标准库。我们将跟随一位开发者的实践经验，了解测试方法、性能对比以及当前面临的挑战。

背景介绍

我的名字是 William。我所在的团队负责为嵌入式 ARM 平台开发工具链，这包括编译器和一些标准库实现。

目前，我们的工具链配置包含一个第三方依赖项 beipy。理想情况下，我们希望用 LLVM-libc 替换它，以实现从源码树内部构建。

但在进行替换之前，我们必须回答一个核心问题：LLVM-libc 是否正确且性能足够？

测试正确性 ✅

为了验证正确性，我们在一个名为 Ar tool chain 的下游仓库中创建了如下测试环境。

这个环境与本地运行测试的常规设置非常相似，关键区别在于我们需要使用模拟器来运行测试套件。同时，还需要链接一些裸机平台特定的文件（图中已用黄色高亮显示）。

运行这些测试后，大部分测试都能通过。目前仅有少数浮点运算测试失败，我们正在努力修复这些问题。

因此，我们现在可以确认 LLVM-libc 在功能上基本是正确的。

评估性能 📊

确认正确性后，我们需要确保其性能达标。为此，我们建立了如下所示的性能对比环境。

请注意，图中用灰色高亮的 C 库意味着我们可以在 Picolypsy 和 LLVM-libc 之间进行切换。

当我们运行性能对比时，在某些基准测试中观察到了性能提升。然而，我们也发现了一些性能回退。

代码大小分析 📏

代码大小对嵌入式应用至关重要。在我们的评估中，我们发现代码大小增加了约 70%。

这种增长主要是由 f theft 函数导致的。目前，Google 的一些开发者正在努力优化代码大小。

核心结论

那么，回到最初的问题：我们能否使用 LLVM-libc 作为标准库？

答案是目前还不能。

虽然已经取得了进展，但仍有许多工作要做。我们相信未来能够实现这个目标。

总结

本节课中，我们一起学习了测试和评估 LLVM-libc 的流程。我们了解了如何搭建测试环境来验证其正确性，以及如何通过性能基准测试和代码大小分析来评估其是否适用于嵌入式场景。尽管 LLVM-libc 已展现出潜力，但在代码大小等方面仍需进一步优化才能完全满足嵌入式应用的要求。

065：迈向 Triton 内核的精确调试 🐛

概述

在本教程中，我们将学习一个名为 TritonSan 的工具，它旨在帮助开发者精确调试基于 Python 的 GPU 编程语言 Triton 中的程序错误。我们将了解 TritonSan 如何利用成熟的 LLVM Sanitizer 技术，通过将内核执行卸载到 CPU 上来检测缓冲区溢出、未初始化内存使用和数据竞争等难以发现的错误。

Triton 语言简介 🚀

Triton 是一种用于高性能计算和人工智能领域的、基于 Python 的 GPU 编程语言。它以学习曲线平缓和语法简洁而闻名，同时又能表达出高效的内核代码。

与任何编程语言一样，如果错误地使用其提供的结构，可能会在代码中引入一些难以察觉的程序错误。目前，可用于帮助程序员检测这些隐蔽错误的工具非常有限。

以下是三类在 Triton 代码中难以检测的编程错误，也是 TritonSan 工具重点关注的类型：

  • 缓冲区溢出
  • 未初始化内存的使用
  • 数据竞争

这些问题之所以难以检测，是因为它们通常在代码中表现出静默或非确定性的行为。例如，在 GPU 上，程序有时会被分配远超其实际需要的内存。如果缓冲区溢出发生在程序分配的总内存范围内，它可能不会被标记为非法访问，从而静默地发生，导致代码出现意外行为。

Triton 的工作流程与错误示例 ⚙️

上一节我们介绍了 Triton 语言及其面临的调试挑战。本节中，我们来看看 Triton 的即时编译工作流程，并通过具体代码示例理解错误是如何发生的。

Triton 使用即时编译。Python 程序逐行执行，直到遇到 GPU 内核调用。此时，该内核会被即时编译。这个过程经过一系列编译器传递：Triton 前端、MLIR 传递，最后通过 GPU 特定的后端编译成可在目标 GPU 上运行的可执行文件。

以下是一个包含数据竞争错误的 Triton 程序示例：

@triton.jit
def kernel(X):
    pid = tl.program_id(axis=0)
    block_start = pid * 2
    offsets = block_start + tl.arange(0, 2)
    mask = offsets < X.shape[0]
    x = tl.load(X + offsets, mask=mask)
    tl.store(X + offsets, x + 1.0, mask=mask) # 数据竞争发生在此行

在这个例子中，我们定义了一个大小为 256 的张量和 128 个同时向该张量写入数据的程序实例。由于程序员错误地将 block_start 设置为 0，导致所有实例都试图同时写入索引 0 和 1，从而引发了数据竞争。

以下是另一个缓冲区溢出的例子：

@triton.jit
def kernel(X):
    pid = tl.program_id(axis=0)
    block_start = pid * 2
    offsets = block_start + tl.arange(0, 2)
    mask = True  # 错误：掩码始终为真
    x = tl.load(X + offsets, mask=mask)
    tl.store(X + offsets, x + 1.0, mask=mask) # 缓冲区溢出发生在此行

这里定义了一个大小为 3 的张量和 2 个写入实例。第二个实例试图写入索引 3，但该索引不存在。错误在于程序员将掩码设置为始终为 True，使得每个程序都认为其访问在范围内，从而导致了缓冲区溢出。

这些小型程序被称为“微基准测试”，用于评估 TritonSan 等工具的精度。

TritonSan 的工作原理 🛠️

上一节我们通过示例了解了 Triton 中典型的错误。本节中，我们来看看 TritonSan 工具是如何工作的。

TritonSan 的核心思想是利用现有的 LLVM Sanitizer。我们选择这样做是因为这些 Sanitizer 成熟、可靠且持续维护。为了利用它们，我们需要将内核执行卸载到 CPU 上，因为这些 Sanitizer 是为在 CPU 上运行而构建的。这样做也带来了优势，例如能够利用 CPU 编译中内置的更成熟、更健壮的 MLIR 和 LLVM 传递，并能使用 OpenMP 等现有 LLVM 运行时来帮助并行化程序。

以下是 TritonSan 工作流程的详细步骤：

  1. 前端处理：Triton 内核首先被 Triton 前端摄取，并转换为 Triton 方言。
  2. 自定义 MLIR 传递：引入一个自定义的 MLIR 传递，添加一些线性代数操作。
  3. 降低至 LLVM IR：利用现有的 LLVM 和 MLIR 内置传递，将代码一直降低到 LLVM IR。
  4. Sanitizer 插桩：在此阶段引入 Sanitizer 标志，将代码编译为目标文件。
  5. 驱动脚本与链接：该目标文件与一个 C++ 驱动脚本结合。该脚本负责执行 Triton 网格的每次迭代。所有这些被组合成一个共享库。
  6. 加载与执行：通过 Python 的 importlib 库立即在 CPU 上加载并启动该共享库。

为了使工具正常工作，我们还需要进行以下关键调整：

  • 并行化执行：默认情况下，驱动脚本按顺序运行 Triton 网格的每次迭代，这意味着不会发生数据竞争。我们使用 OpenMP 并行化内核执行，为每次迭代分配自己的工作线程，从而允许它们并行执行并暴露数据竞争。
  • 添加调试信息：最初报告的错误信息对程序员帮助不大。我们添加了一个 MLIR 传递，以 DI 属性的形式将调试信息插入 IR，并确保这些信息在后续的所有 LLVM 和 MLIR 传递中得以保留，从而在运行时能显示包含文件名和行号的有用错误信息。
  • 预加载 Sanitizer：确保在 Triton 程序开始执行之前预加载 Sanitizer，使其有机会初始化并分配自己的“影子内存”来跟踪执行元数据。
  • 注册钩子回调：拦截操作系统级别的 API 调用，如内存分配或访问。
  • 使用抑制列表：告诉 Sanitizer 忽略哪些非 Triton 模块的错误，以减少无关的误报，专注于内核内的重要错误。

工具评估：精度与性能 📊

上一节我们深入探讨了 TritonSan 的内部机制。本节中，我们通过评估来看看它的实际效果。

我们使用两个主要数据集来评估 TritonSan：

  1. 微基准测试套件：包含特定错误的小型程序。
  2. 注入错误的现实案例：来自官方 Triton 教程和 Triton-Bench 仓库的更复杂内核。

我们从精度和性能两个方面进行评估：

  • 精度：指工具是否能检测到数据集中我们感兴趣的 bug。
  • 性能：指工具引入的运行时的内存和时间开销是否合理。

在精度方面，TritonSan 能够在微基准测试套件中检测到所有我们感兴趣的问题，且没有产生误报。

为了对比，我们还评估了另一个工具 Compute Sanitizer 的精度。Compute Sanitizer 来自 NVIDIA，是 Triton 官方推荐的、用于检测在 NVIDIA GPU 上运行的 Triton 内核中隐蔽程序错误的解决方案。

以下是精度对比结果：

  • 缓冲区溢出：TritonSan 能够检测到所有溢出。Compute Sanitizer 只能检测到“大”的溢出，这可能与 GPU 内存分配机制有关。
  • 数据竞争：TritonSan 能够检测到全局内存中的数据竞争。Compute Sanitizer 不支持检测共享内存中的数据竞争，而 Triton 的同步结构对程序员是透明的，因此数据竞争可能发生在全局内存中。
  • 未初始化内存使用：TritonSan 尚未集成对应的 LLVM Sanitizer，这是未来的开发步骤。

在性能方面，TritonSan 目前支持的 AddressSanitizer 和 ThreadSanitizer 的性能数据与其他应用场景相似，我们认为这对于工具的可用性来说是一个好迹象。

安装与使用指南 📖

上一节我们评估了 TritonSan 的效果。本节中，我们来看看如何安装和使用它。

TritonSan 已作为一个可选功能集成到名为 Triton-Shared 的 CPU 后端中。因此，你需要克隆 Triton-Shared 仓库。该仓库中提供了一个现成的构建脚本，它会安装兼容版本的 LLVM、Triton 和 Triton-Shared，并生成一个便于运行时调用的脚本。

以下是安装和使用步骤：

  1. 安装 TritonSan：

    git clone <triton-shared-repo-url>
cd triton-shared
./build.sh  # 运行构建脚本

  2. 修改 Triton 程序以使用 CPU 后端：
    在你的 Triton Python 文件顶部添加以下两行，告诉 Triton 使用 CPU 后端编译代码：

    import os
os.environ[‘TRITON_INTERPRET’] = ‘1’

    由于在 CPU 上运行内核，你还需要确保张量没有卸载到 GPU。

  3. 使用 TritonSan 运行程序：
    使用 triton-san 驱动脚本来调用 Sanitizer。指定 Sanitizer 类型（目前支持 asan 用于缓冲区溢出，tsan 用于数据竞争），然后加上你通常用来运行 Triton 程序的命令。

    ./triton-san asan python your_program.py  # 检测缓冲区溢出
./triton-san tsan python your_program.py  # 检测数据竞争

总结与问答要点 🎯

在本教程中，我们一起学习了 TritonSan 工具，它通过结合 LLVM Sanitizer 和 Triton 的 CPU 后端，为检测 Triton 内核中隐蔽的程序错误提供了更精确的解决方案。

核心要点总结：

  • 动机：GPU 编程语言 Triton 中难以调试的隐蔽错误（如缓冲区溢出、数据竞争）缺乏有效的检测工具。
  • 解决方案：TritonSan 将内核执行卸载到 CPU，利用成熟的 LLVM Sanitizer（AddressSanitizer, ThreadSanitizer）进行动态分析。
  • 关键技术：通过 OpenMP 实现并行化以暴露数据竞争；添加调试信息以精确定位错误；使用抑制列表减少误报。
  • 评估结果：在检测精度上优于现有的 Compute Sanitizer（特别是在缓冲区溢出和数据竞争方面）；性能开销在可接受范围内。
  • 使用方式：作为 Triton-Shared CPU 后端的一个可选功能，通过简单的环境变量和脚本进行调用。

问答环节要点回顾：

  1. 局限性：TritonSan 要求内核能被其 CPU 后端编译。某些仅限 GPU 运行或不被该后端支持的内核无法使用此工具。
  2. CPU vs GPU Sanitizer：在 CPU 上运行 Sanitizer 的主要优势是能直接利用 LLVM 生态中成熟、健壮的工具链。虽然硬件架构不同，但对于错误检测（而非性能调优）这一主要目的，CPU 环境通常足够。
  3. 精度差异原因：与 Compute Sanitizer 的精度差异可能源于两者不同的内存分配和访问检查机制（如 CPU 的 malloc 与 GPU 的内存分配）。
  4. 性能：对于受支持的内核，在 CPU 上运行 Sanitizer 虽然比原生 GPU 执行慢，但对于调试目的通常是可接受的。

TritonSan 扩展了 LLVM Sanitizer 在基于 Python 的 DSL 中的应用，为 Triton 开发者提供了一个有价值的调试工具补充。未来的工作可能包括集成未初始化内存检测的 Sanitizer，以及进一步扩展 CPU 后端以支持更多内核类型。

066：迈向面向集合的编译

在本节课程中，我们将探讨如何在LLVM中迈向面向集合的编译。我们将了解当前LLVM在处理数据集合时的局限性，并介绍一个名为Memoir的解决方案，它通过扩展LLVM IR，将数据集合提升为一等公民，从而为分析和优化开辟新的可能性。

当前LLVM中集合表示的局限性

上一节我们介绍了LLVM作为多种编程语言的共享中间表示。本节中我们来看看数据集合这一常见特性。

数据集合为组织数据组、执行操作、操作数据及访问元素提供了逻辑表示。它们蕴含了丰富的语义信息，可用于分析和优化。例如，它们分离了存储在集合内的数据与集合本身的组织结构。在某些语言中，它们还提供了额外的类型保证。

然而，当我们审视LLVM时，它对集合的视图是较低层次的。具体来说，集合被降低为内存块和指向这些内存块的指针。因此，我们的数据和描述集合组织的元数据都必须被降低到这些内存块中，导致数据和元数据混杂在一起。

结果就是，当我们查看程序中的加载或存储指令时，无法分辨我们正在查看的是数据还是集合本身的元数据，这使得分析这些程序变得非常困难。此外，在很大程度上，这些类型保证基本被丢弃了，我们只剩下一些额外的元数据信息。

缺乏集合表示带来的问题

正如之前提到的，这种表示的缺失给我们的分析和优化带来了问题。让我们来看一个具体的例子。

以下是C++中使用无序映射的一个示例：

std::unordered_map<int, int> map;
map.insert({1, 10});
map.insert({2, 20});
std::cout << map[1]; // 期望总是输出10

人类很聪明，知道打印语句总是会输出常量值10。但我们正在构建的编译器并没有那么智能。目前，没有任何生产级编译器能够将常量10传播到下面的打印语句中。

这是因为第二次插入操作可能导致集合的重新哈希。结果，键1对应的逻辑元素可能被移动到一个新的内存位置。由于我们的编译器只在内存位置的概念上操作，我们丢失了关于逻辑元素及其值未改变的所有信息。

Memoir：面向集合的LLVM IR扩展

为了解决这类问题，我们引入了Memoir。它通过将数据集合作为SSA形式中的一等公民，扩展了LLVM IR。

Memoir接收来自C、C++或Rust的代码，并生成扩展的LLVM IR。这种扩展包括为这些集合添加额外的类型，以及用于访问和更新它们的操作符。

以下是Memoir中集合类型和操作的示意：

// 伪代码表示Memoir扩展
%map = memoir.collection.alloc unordered_map<int, int>
memoir.collection.insert %map, {1, 10}
memoir.collection.insert %map, {2, 20}
%val = memoir.collection.access %map, 1

Memoir带来的优化能力

有了Memoir，我们能够执行以集合为中心的优化。以前面的代码为例，现在它被翻译成了Memoir的SSA形式。利用Memoir，我们可以执行之前所需的静态分析。这种分析既易于实现，成本又低，因为它本质上是一种数据流分析。

利用分析结果，我们能够执行之前期望的优化，现在打印语句总是输出常量10。我们还可以消除现在已失效的读取操作。

除了这类简单优化，Memoir还能在这些程序上启用一种新的数据转换范式。例如：

  • 安全地执行死字段消除：移除集合中不再使用的数据字段。
  • 缓存局部性优化：例如将“冷”字段从“热”对象中迁移出去。
  • 基于用例的内存布局特化：例如，根据程序使用模式，使用位集（bitset）替代哈希集合（hashset）。

Memoir的性能与内存优势

通过Memoir，我们能够提供显著的性能提升，其幅度远大于当前使用LLVM、GCC和ICC等工具所能达到的水平。此外，Memoir能够减少程序的内存使用量，这是现代编译器目前不太具备的新能力。

优化完成后，我们将Memoir集合降低为来自C++标准库、Abseil或Boost的具体实现。这会产生一个LLVM程序，然后我们将其传递给现有的通用编译流水线。

Memoir的易用性与未来计划

目前，我们正在努力使Memoir更易于使用。具体来说，我们现在有一个C/C++库，你可以用它向编译器暴露这些集合以进行优化。

我们正致力于使其成为一个即插即用的STL替代品，以便更容易地集成到你现有的代码库中。未来，我们计划通过一个Clang前端插件使其完全透明化。然而，目前我们正在进行一些工作，开发一个AI移植工具。如果你等不及完全透明化的方案，这个工具可以将你当前的STL代码转换为使用Memoir集合。

总结

本节课中，我们一起学习了当前LLVM在处理高级数据集合语义时的挑战。我们探讨了Memoir如何通过扩展LLVM IR，将集合作为一等公民引入，从而保留丰富的语义信息，并实现更强大、更高效的编译器分析和优化。Memoir不仅提升了性能，还降低了内存占用，并为未来的编译器优化开辟了新的道路。

067：走近LLVM基金会

概述

在本节课中，我们将了解LLVM基金会的基本情况、使命、运作方式以及当前面临的挑战与机遇。我们将学习基金会如何支持LLVM社区，并探讨其财务现状与未来计划。

基金会简介

我的名字是Tanya Lattner，我是LLVM基金会的执行董事兼主席。欢迎参加这次非正式的讨论，我们将介绍基金会项目的一些更新，并回答大家的问题。

LLVM基金会是一个501(c)(3)公共慈善组织，成立于2014年。作为一个公共慈善机构，我们服务于公共利益。我们有一个由11人组成的董事会，每两年选举一次，目前基金会有两名员工。

我们的使命如下，其中最后一段最为重要：我们通过帮助社区成长、促进社区互动、通过基础设施保持LLVM开发的高效性，并努力确保LLVM项目的长期健康，来支持LLVM社区。我们希望LLVM在未来许多年都能持续成长和繁荣。

基金会项目

我们通过以下项目来执行我们的使命：

  • 教育推广：包括教育材料和活动，例如LLVM开发者大会。
  • 社区0：这是我们的多元化和包容性推广计划，提供奖学金和资助。
  • 社区健康与成长：支持项目基础设施、法律问题以及其他跨领域的工作。

董事会成员介绍

现在，我想借此机会介绍一下LLVM基金会的董事会成员。

  • Chris Bieneman：我是董事会财务主管，同时也是微软的工程师。
  • Chris Lattner：你们可能知道我做的某些事情，我主要在董事会提供建议或帮助解决争议。
  • Mike Edwards：我自2018年起担任董事会成员，之前负责财务工作，现在主要在财务委员会帮忙。
  • Anna：我在苹果工作，一年前加入董事会，但参与LLVM社区已超过18年。我非常乐意帮助新人，并为未来几代人服务。
  • Kristof Beyls：我加入董事会几年了，负责一些事务。
  • An：我是一名LLVM贡献者，也在帮助董事会处理各种事务。你们可能通过Google Summer of Code认识我。
  • David Chisnall：我是董事会的新成员，从LLVM还是学校项目时就开始参与。我热爱社区的构建方式以及基金会的帮助。
  • Hera：我在剑桥大学从事硬件安全工作，也是Linux对冷门安全硬件架构的支持者。我是董事会新成员，但曾在其他八个开源基金会董事会任职，帮助提供历史知识和与其他基金会的联系，并帮助支持基础设施团队获取CI资源。

此外，有两位董事会成员今天无法到场：Wei Wu和Ralf K?chners（董事会秘书）。

加入董事会

如果你有兴趣参与，可以加入我们的董事会。我们是自我延续的董事会，每两年进行一次选举，下一轮选举将在2026年8月进行。

我们正在寻找具备不同技能的人来帮助我们完善董事会并实现目标。所需的技能包括：

  • 非营利组织经验
  • 法律经验
  • 会计和财务经验
  • 先前董事会服务或开源倡导、社区建设等经验

这并非详尽列表，只是一个示例。但最重要的是，你必须拥有帮助LLVM繁荣发展的热情。

我们有一个申请流程，包括与董事会成员的面试。我们可能会在未来几个月内稍微修改这个流程，并会就此进行沟通。如果你有兴趣了解更多关于日常工作或月度工作的情况，可以与在场的任何董事会成员交谈。

项目更新：教育推广

接下来，我们进入项目更新部分，主要讨论教育推广项目。如果你对其他项目有疑问，可以在最后提出。

我们在教育推广项目上取得了巨大成功，本次会议就是明证。我们每年持续组织两次开发者大会。今年六月，我们在亚洲地区（东京）举办了首次活动，有130名与会者。我们看到活动持续增长，几乎回到了疫情前的水平。所有的演讲都会被录制并发布，这是关键的知识库。

然而，我们也面临挑战，因为成本非常高且持续上涨。

这张图表显示了多年来欧洲和美国LLVM开发者大会的出席人数。你可以忽略2020年至2022年的缺口，因为疫情期间我们没有持续举办活动。但如图所示，我们开始接近2019年的水平。

另一方面，正如我们提到的，成本正在上升。从2018年到2024/2025年，这些活动的成本增加了一倍多，这是一个巨大的增长。

尽管如此，得益于赞助商、个人和企业支持者的慷慨捐赠，我们仍然努力使活动尽可能成功。我们仍然能够逐年补贴门票价格。

这张图表显示了开发者大会每年的收入和支出，这是另一种查看方式。你可以看到，我们的门票收入并不完全能覆盖活动的实际支出。

财务状况

如果你对我们除了教育推广或活动之外的资金去向感兴趣，开发者大会确实占我们支出的绝大部分，你可以在此饼图中看到亮蓝色的部分。第二大支出是薪酬。奖学金、资助、社区0和基础设施等其他项目占比较小。这些图表展示了我们2024年的支出以及2025年的预测。

让我们谈谈财务状况。从2017年到2024年，我们的赞助收入增长了14%。然而，同期我们的支出增长了66%。简而言之，维持我们现有运营的成本越来越高。

我们的赞助计划在过去10年里变化很小。因此，鉴于活动成本上升以及我们实现使命的能力，董事会成立了一个专门的子委员会，专注于改革我们的赞助计划，以帮助我们实现一些财务目标。

这是我的呼吁：我们确实在寻找更多赞助商，以帮助我们弥合支出与收入之间的差距，同时继续使我们的活动易于参与。因此，如果你有兴趣赞助LLVM基金会，请联系我或任何董事会成员。

问答环节

现在，我想深入探讨你们的问题。我将开放提问，任何董事会成员如果想对财务讨论做出更多贡献，也可以发言。

提问：关于门票价格与赞助
（Chris Lattner补充）我想指出这张图表。蓝线是我们的赞助收入，红线是我们的总收入，黄线是我们的总支出。目前，我们主要通过提高门票价格来平衡预算。我们一直在讨论的是，门票价格现在高达1200美元，这实际上已经到了一个地步，即使是公司的员工，公司也可能因为价格昂贵而决定减少派遣人数。我们正在讨论是否应该将门票价格降低，比如降到500美元以下。这张图表清楚地表明，如果我们想做到这一点，我们几乎需要将赞助翻倍。因此，我们未来12个月的重点之一将是：如何引入更多资金？如何实现更可持续的基础？另外，图表未显示但很重要的一点是，为了举办这次活动，我们在收到任何门票收入之前，就必须预先支付大约30万美元。我们必须维持足够的现金储备来覆盖所有这些预付款。目前我们的现金储备尚可，但随着活动成本增加，五年后可能就不够了。因此，我们也在思考如何确保拥有可持续的基础和收入增长，以在未来支持社区。

提问：关于参会人数上限
我们总是基于场地空间和物流对参会人数有一定限制。此外，因为我们向酒店保证了一定的消费额，并且我们尽量保守以避免超额支付，所以也会有一些限制。今年我们没有售罄，最多可容纳550人。尽管今年没有售罄，但我们的参会人数同比增长了近13%，增长非常显著。

提问：最大的成本是什么？
最大的成本是餐饮，而不是场地租金。

提问：关于基金会法律结构的影响
（Chris Lattner回答）LLVM基金会是501(c)(3)非营利组织，这与教堂等属于同一类别。这意味着我们的章程是确保社区成功。其他许多非营利基金会，如Linux基金会，是501(c)(6)组织，这类组织通常是会员制，可能导致“付费即玩”的情况，公司可以通过投入资金获得权利等。这导致公司可能为Rust（注：此处为举例，Rust基金会也是501(c)(6)）投入比LLVM更多的资金，这看起来不太合理。我们选择(c)(3)结构是为了确保LLVM的长期可持续性，并有一个由选举产生的董事会来确保LLVM长期成功，而不是被公司政治左右。但另一方面，这带来了资金挑战。因此，我们需要帮助向你们的组织解释基金会的价值，沟通和理解这一点非常重要。

提问：关于税收抵扣
（Mike Edwards补充）作为501(c)(3)组织，赞助商给我们的捐款赞助费是完全免税的。这是一个很好的卖点，可以告诉你的老板，如果他们开出更大的支票，年底报税时可以获得相应的更大税收减免。

提问：是什么阻止了大公司赞助商捐赠更多钱？
我认为关键是我们从未要求过更多资金。我们的赞助商年复一年非常慷慨地开出支票，但可能10年前批准的金额就一直沿用至今，因为一旦进入支付流程，再去要求更改很困难。未来一年，我们的财务委员会和赞助委员会将会主动去敲门，礼貌地请求更多帮助，因为做所有这些事情的成本越来越高了。

（Chris Bieneman补充）董事会需要做的是，清楚地告诉赞助商我们将用这些钱做什么，以及这将如何为他们的员工和整个生态系统带来价值。我们需要更好地阐明这一点，这是我们今年要做的事情。

提问：关于其他收入来源（如拨款、个人捐赠）
我们目前没有考虑拨款，过去也没有。这不意味着我们不会改变主意，但这取决于拨款是否附带条件。至于个人捐赠者，你可以在LLVM网页上找到捐赠按钮直接捐款。我们也尝试通过商品等方式激励捐赠，但这部分收入很少。此外，今年我们加入了Github Sponsors，你可以在Github上赞助我们。

（David Chisnall补充）个人捐赠是免税的。许多人可能不知道，如果你捐赠增值资产（如股票），可以避免资本利得税，同时获得慈善捐赠的税收减免。你可以将股票捐赠给基金会，基金会出售时无需支付资本利得税，而你获得全额捐赠价值的税收减免。此外，你的公司可能提供企业匹配捐赠，你应该查看是否适用于向LLVM基金会的捐赠。

（Chris Lattner补充）捐赠增值股票可以获得三到五倍的捐赠价值效益，如果再结合公司匹配，效益倍数更高。这对慈善机构非常有益。

提问：关于增加赞助商数量
我们绝对没有达到上限。有很多使用LLVM的公司目前并未赞助LLVM基金会。作为董事会，我们需要改进对公司的外联工作，提出赞助请求。我们之前要求得不够多，这可能是因为很长时间以来我们不需要这样做。此外，今年在亚洲举办活动也是为了开拓新市场，开辟更多赞助机会。建立赞助渠道需要时间，我们有很多工作要做。使用或基于LLVM构建产品的公司数量巨大，其中很多不在我们的赞助商名单上，因此我们有很多机会。

提问：LLVM的贡献是否集中在少数几家公司？
不，绝对在增长。我们看到每年都有新公司加入这个项目。LLVM项目正在广泛增长，社区变得如此庞大和活跃，我们需要思考如何将这种增长也带入基金会的赞助中。

总结

本节课中，我们一起学习了LLVM基金会的使命、组织结构和核心项目。我们了解了基金会主要通过教育推广（如开发者大会）、社区0和社区健康项目来支持LLVM生态。同时，我们也深入探讨了基金会当前面临的财务挑战，特别是活动成本上升与收入增长之间的差距，以及董事会为寻求更可持续的运营模式所做的努力。最后，我们看到了社区参与和支持（无论是通过赞助、个人捐赠还是加入董事会）对于基金会和LLVM项目长期健康发展的重要性。

感谢大家的参与。请随时与任何董事会成员交流，提出更多问题或分享你对社区的需求和想法。本次会议结束后，大家可以前往海报展示区，那里有一些点心和咖啡。

068：数据流分析中实际传递函数的合成

概述

在本教程中，我们将学习如何在数据流分析中合成实际的传递函数。我们将从数据流分析的基础概念开始，逐步解释已知位分析，探讨手动编写传递函数的挑战，并最终介绍一种能够自动合成正确且精确的传递函数的方法。

数据流分析与已知位分析回顾

上一节我们概述了本教程的目标，本节我们来回顾数据流分析的基本概念。数据流分析旨在获取在所有程序执行路径上都成立的信息或属性。

右侧是一个简单的LLVM IR示例函数。数据流分析试图获取在所有执行中都成立的信息。一个具体的例子是“已知位”分析。它试图确定某个比特位在所有执行中是否总是0或总是1。

对于已知位，有三种可能性：

  • 0：在所有执行中，该位都是0。
  • 1：在所有执行中，该位都是1。
  • 未知：该位可能是0，也可能是1。

根据这个定义，我们来看第一个操作 %n3 = and i4 %a, 3。%a 是一个4位宽的未知操作数，3 是常量。经过这个 and 操作后，我们知道结果的高两位总是0，低两位未知。这意味着 %n3 可能的运行时值是0、1、2、3。

类似地，对于下一个操作 %n1 = and i4 %b, 1，我们知道结果的高三位总是0，最低位是1。因此 %n1 可能的运行时值是0和1。这就是已知位信息。

LLVM使用两个成员变量来实现已知位信息：KnownZero 和 KnownOne。右侧附有两个示例：

  • 如果一个位总是0，那么 KnownZero 中对应的位被置位。
  • 如果一个位总是1，那么 KnownOne 中对应的位被置位。
  • 如果未知，则两个位都不置位。

传递函数的作用

既然我们知道了 %n3 和 %n1 的已知位信息，如何推断出 %xor 结果的已知位信息呢？

右侧附有异或操作的已知位真值表。通过应用这个真值表，我们可以得到结果。在LLVM中，他们通过一个传递函数来实现这个真值表。这个函数接收两个操作数（%n3 和 %n1）的已知位信息，执行几行代码来计算，并将结果信息附加到 %xor 上。

在这个例子中，传递函数接收两个操作数的已知位信息，并返回结果的已知位信息。这是一个简单的异或操作转换。

LLVM为已知位域中的不同操作实现了许多传递函数，其中一些非常复杂。例如，加法操作的传递函数 computeForAddCarry 是一个用于加法和减法的实用函数。它接收两个已知位操作数，调用 getMinValue、getMaxValue 等函数，结合位运算，最终返回结果。这个过程至少可以说不够直观。

总结一下，已知位只是LLVM提供的一个分析域。还有其他域，如整数范围分析或活跃变量分析。LLVM为不同域上的不同操作实现了各种传递函数。

为什么需要合成器？

上一节我们了解了传递函数的作用，本节我们来看看为什么需要自动合成器。首先，让我们看看合成器能做什么。

简而言之，合成器读取一个规范。以已知位为例，规范包括：

  1. 已知位域的定义。
  2. 要应用的操作（例如，xor、add）。
  3. 传递函数的签名（例如，接收两个已知位信息作为参数）。

合成之后，我们可以产生一个结果：一个正确的传递函数。它能确保：

  • 正确性：总是给出正确的信息。
  • 合理性：至少返回合理的结果。
  • 可用性：可用于优化或过程间分析。
  • 可嵌入性：可以直接包含在您的项目中，并在数据流分析中调用。

那么，为什么需要合成器呢？首先，当然是为了获得更多的传递函数。

以下是不同平台上已实现传递函数的内在函数数量与内在函数总数的对比表（基于公开信息）：

  • X86 (LLVM)：已实现 30 个，总数 1714 个。
  • ARM：已实现 5 个。
  • AArch64：已实现 2 个。

通过使用合成器，我们可以获得更多传递函数。

另一个例子是MIR（Machine IR）中的CodeGen。CodeGen是LLVM后端项目中使用的一个域，它只包含7个操作的传递函数，而操作总数是20个。对于RISC-V，则完全没有已知位分析。因此，我们相信其他领域（如WebAssembly、AMDGPU）也能从合成器中受益，从而获得更多的传递函数。

您可能会问：没有传递函数也没关系，我可以复用LLVM IR中现有的那些。

接下来，我将解释复用现有传递函数的困难所在。

首先，是的，他们确实尝试过复用。上面的例子是CodeGen中使用的最长的传递函数。如果传入的操作是常量操作，CodeGen会直接复用LLVM的已知位结构，调用 KnownBits::makeConstant 函数并返回结果。

第一个问题是操作语义可能不同。底部有一个关于移位操作的例子。在CodeGen中，如果移位量大于或等于位宽，结果是0。例如，移位量是5，位宽是4，结果是0。

然而，在LLVM IR中，它定义移位量必须小于位宽，否则结果返回 poison。这是第一个问题。

第二个问题涉及MIR中的另一个分析：整数范围分析。它试图确定一个变量是否总是落在某个整数范围内（通过上下界表示）。有一个文件 inferIntRange 提供了一组实用函数，被两个后端（AArch64和AMDGPU）共享。

在顶部的例子中，AArch64后端的 and 操作直接调用了 inferIntRange 文件中的 inferAnd 函数，这很好。

然而，对于AMDGPU后端，它首先调用 inferAnd，然后调用一个适配器函数 inferIndexOp 将结果转换到索引域。这意味着程序员必须手动编写这些适配器，这很容易出错，并且可能包含错误。这是一个挑战。

合成器的最后一个好处是我们可以证明传递函数的正确性。

为了定义正确性，这里引入两个概念：

  1. 正确性：结果必须覆盖所有可能的运行时值。
    • 示例：假设实际的运行时值集合是 {0}。如果一个错误的不正确结果说“未知”（可能的值是 {0, 1}），这仍然是正确的，因为它覆盖了实际值。如果一个错误的不正确结果说“总是1”（可能的值是 {1}），这就不正确，因为它没有覆盖实际值0，可能导致错误的优化。
  2. 精确性：结果中包含的、永远不会在运行时出现的值越少越好。
    • 示例：假设实际的运行时值集合是 {0}。
      • 最精确的结果是“总是0”（{0}）。
      • 一个不那么精确但仍然是正确的结果是“未知”（{0, 1}），因为它包含了不会出现的值1。
    • 在图表中，如果A和B都是正确的，且A是B的子集，那么A比B更精确。

合成器的设计

上一节我们探讨了需要合成器的原因，本节我们来看看合成器的设计。合成器的大致流程是：读取规范，生成一批候选传递函数，评估这些候选函数，将正确的候选函数加入解决方案集，不断修剪解决方案集，最终快速生成解决方案。

现在让我们看看规范包含什么。规范包含三个部分：

  1. 域的定义：例如，对于已知位域，包括域成员（KnownZero, KnownOne）。
  2. 域的约束：例如，在已知位中，一个位不能同时是0和1。
  3. 构造函数：例如，已知位的默认构造函数返回“全未知”；meet 操作接收两个已知位信息，返回它们的交集。
  4. 操作的定义：假设操作必须配备SMT语义（例如，可以将 xor、add 转换为SMT表达式）。操作本身可能包含一些约束（例如，移位量必须小于位宽）。
  5. 传递函数的签名：例如，接收左操作数和右操作数的已知位信息。

有了规范之后，我们从函数签名开始，用随机操作填充函数体，并逐步进行变异。

我们使用随机搜索的原因是搜索空间极其庞大，无法穷举所有可能性。以下是一个单步变异的具体例子。代码片段中包含 autogen9, autogen10, autogen11 等变量。有两种变异可能：

  1. 随机选择一个操作，改变其操作数（例如，将 autogen9 & autogen10 改为 autogen9 | autogen10）。
  2. 改变一个操作数（例如，将第一个操作数从 autogen6 改为 autogen9）。

找到可能的候选函数后，我们需要评估它们。评估引擎从底部接收一个候选函数，将其发送到SMT求解器检查正确性。如果正确，则加入解决方案集；否则拒绝。

另一种方式是在低位宽上进行穷举测试，以评估其精确度。如果一个候选函数正确但精确度很低，我们仍然不想要。

合成器保留两种候选函数：

  1. 正确且高精确度的候选函数。我们保留它。
  2. 仅在特定条件下正确（但仍有高精确度）的候选函数。这意味着它们只在一小部分情况下正确，必须通过前提条件来捕获。对于这类候选函数，我们不是合成函数体，而是合成一个 if 前提条件。如果前提条件为真，则执行候选函数体并返回结果。

通过不断将候选函数加入解决方案集，解决方案集会变得越来越大。我们希望移除那些效率较低的函数。这里我们使用一种贪心策略：

  1. 从旧的解决方案集开始（初始为空）。
  2. 基于新的（当前）解决方案集评估所有候选函数。
  3. 选择贡献最大的候选函数（例如，能处理最多之前未覆盖的输入情况）。
  4. 如果其贡献大于0，则将其加入新的解决方案集，并重复此过程。
  5. 如果最大贡献为0，意味着剩余的候选函数无用，我们停止此过程。

最终，我们将得到一个包含一组候选函数的解决方案集，并生成一个简单的传递函数。我们合并所有候选函数或部分解决方案的结果。以下是合成出的 xor 传递函数示例，它包含两个候选函数（部分解决方案），最终通过 meet 操作合并结果。

另一个更复杂的例子是减法操作，它包含16个部分解决方案，需要多次合并才能返回最终结果。这种方法有效的原因是，我们确保每个候选函数都是正确的，因此合并所有结果后仍然是正确的，因为它必须覆盖所有运行时值。

实验结果与总结

最后，我们来看一些实验结果。我们在已知位域上对39个操作进行了合成测试，使用随机API输入进行评估。

左侧表格显示了我们表现良好的操作：

  • and、or、xor：与LLVM的精确度相同。
  • mul（乘法）：我们达到59%的精确度，优于LLVM的53%。
  • shl.sat（饱和左移）：我们达到79.7%的精确度，而LLVM不支持此操作。
  • mul（另一种形式）：我们达到60%的精确度。

右侧表格显示了我们表现不佳的操作：

  • averageFloor：我们只有不到40%的精确度。
  • sadd.sat、ssub.sat、absdiff：我们只有约60%的精确度，而LLVM接近100%。

结果表明，结果对随机因素（如初始程序和随机变量）非常敏感。

我们还在2017年的CPU上使用SPEC基准测试进行了评估。总体而言，我们做得不算差。与GCC相比，我们比较了LLVM实现的已知位分析和我们合成的分析。我们只损失了1.78%的精确度，甚至与MCF项目达到了相同的结果。在其他项目中，我们有大约3-5%的损失。我们在X64上也做了一些不好的工作，我们认为这可能包含了许多我们的合成器处理不佳的指令。

总结一下，我们的目标不是超越LLVM，而是在面对新的数据域、新的分析域或新的操作时，生成正确且精确的传递函数。通过读取规范以及配备了SMT语义的操作，用户可以生成可用于数据流分析的正确传递函数。

本教程到此结束。

posted @ 2026-03-29 09:18  绝不原创的飞龙  阅读(9)  评论(0)    收藏  举报