如何打破CUDA垄断？LLVM奠基人长文解读

如果您希望可以时常见面，欢迎标星收藏哦~

编者按：

多年来，领先的人工智能公司一直坚称，只有拥有庞大计算资源的公司才能推动前沿研究，这强化了这样一种观点：除非你拥有数十亿美元的基础设施投入，否则“不可能赶上”。但 DeepSeek 的成功却讲述了一个不同的故事：新颖的理念可以带来效率上的突破，从而加速人工智能的发展；规模更小但更专注的团队可以挑战行业巨头，甚至创造公平的竞争环境。

我们认为，DeepSeek 的效率突破预示着AI 应用需求的激增。如果 AI 要继续发展，就必须降低总体拥有成本 (TCO) ——通过扩大替代硬件的覆盖范围、最大限度地提高现有系统的效率以及加速软件创新。否则，未来 AI 的效益将面临瓶颈——要么是硬件短缺，要么是开发者难以有效利用现有的各种硬件。

这不仅仅是一个抽象的问题——这是我（指代本文作者Chris Lattner，下同）整个职业生涯都在努力解决的一个挑战。

过去 25 年来，我一直致力于为世界释放计算能力。我创立并领导了LLVM的开发，LLVM 是一项编译器技术，它为 CPU 在编译器技术的新应用领域打开了大门。如今，LLVM 已成为 C++、Rust、Swift 等性能导向型编程语言的基础。它支持几乎所有 iOS 和 Android 应用，以及 Google 和 Meta 等主要互联网服务的基础设施。

这项工作为我在苹果领导的几项关键创新铺平了道路，包括创建OpenCL（一个早期的加速器框架，现已被整个行业广泛采用）、使用LLVM重建苹果的CPU和GPU软件堆栈，以及开发Swift编程语言。这些经历强化了我对共享基础设施的力量、软硬件协同设计的重要性，以及直观、开发者友好的工具如何释放先进硬件的全部潜力的信念。

2017 年，我开始着迷于 AI 的潜力，并加入 Google，领导 TPU 平台的软件开发。当时，硬件已经准备就绪，但软件尚未投入使用。在接下来的两年半时间里，通过团队的共同努力，我们在 Google Cloud 上推出了 TPU，并将其扩展到每秒百亿亿次浮点运算 (ExaFLOPS)，并构建了一个研究平台，促成了Attention Is All You Need和BERT等突破性成果。

然而，这段旅程也揭示了人工智能软件更深层次的问题。尽管 TPU 取得了成功，但它们仍然仅与 PyTorch 等人工智能框架半兼容——谷歌凭借巨大的经济和研究资源克服了这个问题。一个常见的客户问题是：“TPU 能开箱即用地运行任意人工智能模型吗？”真相是？不能——因为我们没有 CUDA，而 CUDA 是人工智能开发的事实标准。

我并非回避解决行业重大问题的人：我最近的工作是创建下一代技术，以适应硬件和加速器的新时代。这包括 MLIR 编译器框架（目前已被整个行业广泛采用的 AI 编译器），以及我们团队在过去 3 年中构建的一些特别的东西——但我们稍后会在合适的时机分享更多相关信息。

由于我的背景和在业界的人脉，我经常被问及计算的未来。如今，无数团队正在硬件领域进行创新（部分原因是NVIDIA 市值飙升），而许多软件团队正在采用 MLIR 来支持新的架构。与此同时，高层领导们也在质疑，为什么尽管投入了大量资金，AI 软件问题仍然悬而未决。挑战并非缺乏动力或资源。那么，为什么这个行业会感到停滞不前呢？

我不认为我们陷入了困境。但我们确实面临着一些棘手的基础性问题。

为了向前发展，我们需要更好地理解行业底层动态。计算是一个技术含量极高的领域，发展迅速，充斥着各种术语、代号和新闻稿，旨在让每一款新产品都听起来具有革命性。许多人试图拨开迷雾，只见树木不见森林，但要真正理解我们的发展方向，我们需要探究其根源——那些将一切联系在一起的基本构件。

首先，我们将以一种简单易懂的方式回答这些关键问题：

• CUDA 到底是什么？

• CUDA 为何如此成功

• ️CUDA 真的好吗？

• 为什么其他硬件制造商难以提供可比的 AI 软件？

• 为什么 Triton、OneAPI 或 OpenCL 等现有技术还没有解决这个问题？

• 作为一个行业，我们该如何向前发展？

我希望本文能够激发有意义的讨论，并提升人们对这些复杂问题的理解。人工智能的快速发展——例如 DeepSeek 最近的突破——提醒我们，软件和算法创新仍然是推动行业发展的动力。对底层硬件的深入理解继续带来 “10 倍” 的突破。

人工智能正以前所未有的速度发展，但仍有诸多潜力有待挖掘。让我们携手突破，挑战固有认知，推动行业发展。让我们一起深入探索！

“CUDA”究竟是什么？

似乎在过去一年里，每个人都开始谈论 CUDA：它是深度学习的支柱，是新型硬件难以与之竞争的原因，也是英伟达护城河与市值飙升的核心。

DeepSeek 的出现让我们有了惊人的发现：它的突破是通过 “绕过” CUDA、直接进入 PTX 层实现的…… 但这究竟意味着什么呢？似乎每个人都想打破这种技术锁定，但在制定计划之前，我们必须先了解自己面临的挑战。

CUDA 在人工智能领域的主导地位不可否认，但大多数人并不完全理解 CUDA 究竟是什么。有人认为它是一种编程语言，有人称它是一个框架。许多人认为它只是 “英伟达用来让 GPU 运行更快的东西”。这些说法并非完全错误，也有很多杰出人士试图解释它，但没有一种能完全涵盖 “CUDA 平台” 的全貌。

CUDA 并非单一事物，它是一个庞大的分层平台，是一系列技术、软件库和底层优化的集合，共同构成了一个大规模的并行计算生态系统。它包括：

• 一种底层并行编程模型，开发者可以用类似 C++ 的语法利用 GPU 的原始计算能力。

• 一套复杂的库和框架，这些中间件支持人工智能等关键垂直应用场景（例如用于 PyTorch 和 TensorFlow 的 cuDNN 库 ）。

• 像 TensorRT-LLM 和 Triton 这样的高级解决方案，它们能在不需要开发者深入了解 CUDA 的情况下，支持人工智能工作负载（例如大语言模型服务）。

而这只是冰山一角。

在本章节中，我们将深入剖析 CUDA 平台的关键层级，探究其发展历程，并解释它为何对如今的人工智能计算如此重要。这为我们系列文章的下一部分内容奠定了基础，届时我们将深入探讨 CUDA 如此成功的原因。提示：这与其说是技术本身的原因，不如说和市场激励因素有很大关系。

让我们开始吧！

CUDA 发展之路：从图形处理到通用计算

在 GPU 成为人工智能和科学计算的强大引擎之前，它们只是图形处理器，是专门用于渲染图像的处理器。早期的 GPU 将图像渲染功能硬编码在硅芯片中，这意味着渲染的每个步骤（变换、光照、光栅化）都是固定的。虽然这些芯片在图形处理方面效率很高，但缺乏灵活性，无法用于其他类型的计算。

2001 年，英伟达推出 GeForce3，这一切发生了改变。GeForce3 是第一款带有可编程着色器的 GPU，这在计算领域是一次重大变革：

• 在此之前：固定功能的 GPU 只能应用预定义的效果。

• 在此之后：开发者可以编写自己的着色器程序，解锁了可编程图形管线。

这一进步伴随着 Shader Model 1.0 的推出，开发者可以编写在 GPU 上执行的小程序，用于顶点和像素处理。英伟达预见到了未来的发展方向：GPU 不仅可以提升图形性能，还能成为可编程的并行计算引擎。

与此同时，研究人员很快就提出了疑问：“如果 GPU 能运行用于图形处理的小程序，那我们能否将其用于非图形任务呢？”

斯坦福大学的 BrookGPU 项目是早期对此进行的重要尝试之一。Brook 引入了一种编程模型，使 CPU 能够将计算任务卸载到 GPU 上，这一关键理念为 CUDA 的诞生奠定了基础。

这一举措具有战略意义且极具变革性。英伟达没有把计算当作一项附带实验，而是将其列为首要任务，将 CUDA 深度融入其硬件、软件和开发者生态系统中。

CUDA 并行编程模型

2006 年，英伟达推出 CUDA（统一计算设备架构），这是首个面向 GPU 的通用编程平台。CUDA 编程模型由两部分组成：“CUDA 编程语言” 和 “英伟达驱动程序”。

CUDA 是一个分层堆栈，需要从驱动程序到内核的深度集成

CUDA 语言源自 C++，并进行了扩展，以直接暴露 GPU 的底层特性，例如 “GPU 线程” 和内存等概念。程序员可以使用该语言定义 “CUDA 内核”，这是一种在 GPU 上运行的独立计算任务。下面是一个非常简单的示例：

CUDA 内核允许程序员定义自定义计算，这些计算可以访问本地资源（如内存），并将 GPU 用作高速并行计算单元。这种语言会被翻译成 “PTX”，PTX 是一种汇编语言，是英伟达 GPU 支持的最低级接口。

但是程序究竟如何在 GPU 上执行代码呢？这就要用到英伟达驱动程序了。它充当 CPU 和 GPU 之间的桥梁，负责处理内存分配、数据传输和内核执行。以下是一个简单示例：

请注意，这些操作都非常底层，充满了繁杂的细节（如指针和 “幻数（magic numbers）”）。如果出现错误，通常会以难以理解的程序崩溃形式提示。此外，CUDA 还暴露了许多英伟达硬件特有的细节，例如 “warp 中的线程数”（这里暂不深入探讨）。

尽管存在这些挑战，但这些组件让整整一代硬核程序员能够利用 GPU 强大的计算能力来解决数值问题。例如，2012 年 AlexNET 点燃了现代深度学习的火种。它之所以能够实现，得益于用于卷积、激活、池化和归一化等人工智能操作的自定义 CUDA 内核，以及 GPU 提供的强大算力。

虽然大多数人听到 “CUDA” 时，通常想到的是 CUDA 语言和驱动程序，但这远不是 CUDA 的全部，它们只是其中的一部分。随着时间的推移，CUDA 平台不断发展，涵盖的内容越来越多，而最初的首字母缩写词（CUDA）已经无法准确描述其全部意义。

高级 CUDA 库：让 GPU 编程更易上手

CUDA 编程模型为通用 GPU 计算打开了大门，功能强大，但它带来了两个挑战：

• CUDA 使用难度较大。

• 更糟糕的是，CUDA 在性能可移植性方面表现不佳。

为第 N 代 GPU 编写的大多数内核在第 N + 1 代 GPU 上仍能 “继续运行”，但性能往往较差，远达不到第 N + 1 代 GPU 的峰值性能，尽管 GPU 的优势就在于高性能。这使得 CUDA 成为专业工程师的有力工具，但对大多数开发者来说，学习门槛较高。这也意味着每次新一代 GPU 推出时（例如现在新出现的 Blackwell 架构），都需要对代码进行大量重写。

随着英伟达的发展，它希望 GPU 对那些在各自领域是专家，但并非 GPU 专家的人也有用。英伟达解决这一问题的方法是开始构建丰富而复杂的闭源高级库，这些库抽象掉了 CUDA 的底层细节，其中包括：

• cuDNN（2014 年推出）—— 加速深度学习（例如卷积、激活函数运算）。

• cuBLAS—— 优化的线性代数例程。

• cuFFT—— 在 GPU 上进行快速傅里叶变换（FFT）。

• 以及许多其他库。

有了这些库，开发者无需编写自定义 GPU 代码就能利用 CUDA 的强大功能，英伟达则承担了为每一代硬件重写这些库的工作。这对英伟达来说是一项巨大的投资，但最终取得了成效。

cuDNN 库在这一过程中尤为重要，它为谷歌的 TensorFlow（2015 年推出）和 Meta 的 PyTorch（2016 年推出）铺平了道路，推动了深度学习框架的兴起。虽然此前也有一些人工智能框架，但这些是首批真正实现规模化应用的框架。现代人工智能框架中包含数千个 CUDA 内核，每个内核都极难编写。随着人工智能研究的爆发式增长，英伟达积极扩展这些库，以涵盖重要的新应用场景。

CUDA 上的 PyTorch 建立在多层依赖关系之上

英伟达对这些强大的 GPU 库的投入，使得全球开发者能够专注于构建像 PyTorch 这样的高级人工智能框架，以及像 HuggingFace 这样的开发者生态系统。他们的下一步是打造开箱即用的完整解决方案，让开发者完全无需了解 CUDA 编程模型。

全面的垂直解决方案助力AI和GenAI快速发展

人工智能的热潮远远超出了研究实验室的范畴，如今它无处不在。从图像生成到聊天机器人，从科学发现到代码助手，生成式人工智能（GenAI）在各个行业蓬勃发展，为该领域带来了大量新应用和开发者。

与此同时，出现了一批新的人工智能开发者，他们有着截然不同的需求。在早期，深度学习需要精通 CUDA、高性能计算（HPC）和底层 GPU 编程的专业工程师。如今，一种新型开发者（通常称为人工智能工程师）在构建和部署人工智能模型时，无需接触底层 GPU 代码。

为了满足这一需求，英伟达不仅提供库，还推出了交钥匙解决方案，将底层的一切细节都抽象掉。这些框架无需开发者深入了解 CUDA，就能让人工智能开发者轻松优化和部署模型。

• Triton Serving—— 一种高性能的人工智能模型服务系统，使团队能够在多个 GPU 和 CPU 上高效运行推理。

• TensorRT—— 一种深度学习推理优化器，可自动调整模型，使其在英伟达硬件上高效运行。

• TensorRT-LLM—— 一种更专业的解决方案，专为大规模大语言模型（LLM）推理而构建。

• 以及许多（众多）其他工具。

NVIDIA 驱动程序和 TensorRT-LLM 之间存在多个层

这些工具完全屏蔽了 CUDA 的底层复杂性，让人工智能工程师能够专注于人工智能模型和应用，而无需关注硬件细节。这些系统提供了强大的支持，推动了人工智能应用的横向扩展。

整体的 “CUDA 平台”

CUDA 常被视为一种编程模型、一组库，甚至仅仅是 “英伟达 GPU 运行人工智能所依赖的东西”。但实际上，CUDA 远不止如此。它是一个统一的品牌，是一个真正庞大的软件集合，也是一个经过高度优化的生态系统，所有这些都与英伟达的硬件深度集成。因此，“CUDA” 这个术语含义模糊，我们更倾向于使用 “CUDA 平台” 这一表述，以明确我们所谈论的更像是 Java 生态系统，甚至是一个操作系统，而不仅仅是一种编程语言和运行时库。

CUDA 的复杂性不断扩大：涵盖驱动程序、语言、库和框架的多层生态系统

从核心来看，CUDA 平台包括：

• 庞大的代码库：经过数十年优化的 GPU 软件，涵盖从矩阵运算到人工智能推理的所有领域。

• 广泛的工具和库生态系统：从用于深度学习的 cuDNN 库到用于推理的 TensorRT，CUDA 涵盖了大量的工作负载。

• 针对硬件优化的性能：每次 CUDA 发布都会针对英伟达最新的 GPU 架构进行深度优化，确保实现顶级效率。

• 专有且不透明：当开发者与 CUDA 的库 API 交互时，底层发生的很多操作都是闭源的，并且与英伟达的生态系统紧密相连。

CUDA 是一套强大且庞大的技术体系，是整个现代 GPU 计算的软件平台基础，其应用甚至超越了人工智能领域。

既然我们已经了解了 “CUDA” 是什么，接下来就需要明白它为何如此成功。提示一下：CUDA 的成功并非真的取决于性能，而是战略、生态系统和发展势头。在下一篇文章中，我们将探讨是什么让英伟达的 CUDA 软件塑造并巩固了现代人工智能时代。

CUDA 如何取得成功？

如果作为一个生态系统，我们希望取得进展，就需要了解 CUDA 软件帝国是如何占据主导地位的。理论上，存在一些替代方案，比如 AMD 的 ROCm、英特尔的 oneAPI、基于 SYCL 的框架等，但在实际应用中，CUDA 仍然是 GPU 计算领域无可争议的王者。

这一切是如何发生的呢？

答案并不仅仅在于卓越的技术，尽管技术确实起到了一定作用。CUDA 是一个开发者平台，它的成功得益于出色的执行、深入的战略投资、连贯性、生态系统锁定，当然，还有一点点运气。

本章节将深入剖析 CUDA 如此成功的原因，探究英伟达的层层战略 —— 从早期对通用并行计算的押注，到与 PyTorch 和 TensorFlow 等人工智能框架的紧密结合。归根结底，CUDA 的主导地位不仅是软件的胜利，更是长期平台思维的典范。

CUDA 的早期发展

构建一个计算平台的关键挑战之一，是吸引开发者学习并投入其中。如果只能针对小众硬件，就很难形成发展势头。在一期精彩的《Acquired》播客节目中，黄仁勋分享了英伟达早期的一个关键战略，即保持 GPU 的跨代兼容性。这使得英伟达能够利用其已广泛普及的游戏 GPU 用户基础，这些 GPU 最初是为运行基于 DirectX 的 PC 游戏而销售的。此外，开发者可以在低价的台式电脑上学习 CUDA，之后再扩展到价格高昂、性能更强大的硬件上。

这在如今看来或许显而易见，但在当时却是一个大胆的举措：英伟达没有为不同的应用场景（笔记本电脑、台式机、物联网、数据中心等）打造单独优化的产品线，而是构建了一条连续统一的 GPU 产品线。这意味着要做出一些权衡，比如在功耗或成本效率方面有所牺牲，但作为回报，它创建了一个统一的生态系统，开发者在 CUDA 上的投入可以从游戏 GPU 无缝扩展到高性能数据中心加速器。这种策略与苹果维护和推动 iPhone 产品线发展的方式颇为相似。

这种方法有两个好处：

• 降低准入门槛：开发者可以使用已有的 GPU 学习 CUDA，便于进行实验和采用。

• 产生网络效应：随着越来越多的开发者开始使用 CUDA，更多的软件和库被开发出来，这让该平台变得更有价值。

这种早期的用户基础使 CUDA 的应用范围从游戏领域扩展到科学计算、金融、人工智能和高性能计算（HPC）领域。一旦 CUDA 在这些领域获得关注，它相较于其他替代方案的优势就变得很明显：英伟达持续的投入确保了 CUDA 始终处于 GPU 性能的前沿，而竞争对手却难以构建一个与之相媲美的生态系统。

抓住并乘上人工智能软件的浪潮

随着深度学习的爆发，CUDA 的主导地位得以巩固。2012 年，开启现代人工智能革命的神经网络 AlexNet，是使用两块英伟达 GeForce GTX 580 GPU 进行训练的。这一突破不仅表明 GPU 在深度学习方面速度更快，还证明了它们对人工智能的发展至关重要，使得 CUDA 迅速成为深度学习的默认计算后端。

随着深度学习框架的出现，尤其是谷歌在 2015 年推出的 TensorFlow 和 Meta 在 2016 年推出的 PyTorch，英伟达抓住机遇，大力投入优化其高级 CUDA 库，以确保这些框架在其硬件上尽可能高效地运行。正如我们在第 2 部分所讨论的，英伟达积极优化 cuDNN 和 TensorRT，而不是让人工智能框架团队自行处理底层的 CUDA 性能调优。

这一举措不仅使 PyTorch 和 TensorFlow 在英伟达 GPU 上的运行速度大幅提升，还让英伟达能够紧密整合其硬件和软件（这一过程被称为 “硬件 / 软件协同设计”），因为这样减少了与谷歌和 Meta 之间的协调成本。英伟达每推出新一代的主要硬件，就会发布一个新版本的 CUDA，以利用新硬件的功能。人工智能领域渴望速度和效率，非常愿意将这项工作交给英伟达，这直接导致这些框架与英伟达硬件紧密绑定。

但谷歌和 Meta 为什么会任由这种情况发生呢？实际上，谷歌和 Meta 并非专注于构建一个广泛的人工智能硬件生态系统，它们更关注利用人工智能来推动营收增长、改进产品以及开展新的研究。它们的顶尖工程师优先处理那些对公司内部指标有重大影响的项目。例如，这些公司决定打造自己的专有 TPU 芯片，将精力投入到为自家的硬件进行优化上。因此，把 GPU 相关的事务交给英伟达处理也就顺理成章了。

替代硬件制造商则面临着艰巨的挑战，他们试图复制英伟达庞大且不断扩展的 CUDA 库生态系统，但却没有像英伟达那样专注于硬件整合。竞争对手不仅进展艰难，还陷入了一个恶性循环，总是在追赶英伟达硬件上的下一个人工智能进展。这也影响了谷歌和 Meta 的内部芯片项目，催生了包括 XLA 和 PyTorch 2 在内的众多项目。我们会在后续文章中深入探讨这些内容，但如今可以看到，尽管人们抱有一些期望，却没有什么能让硬件创新者具备与 CUDA 平台相媲美的能力。

英伟达凭借每一代新硬件不断扩大领先优势。然后，在 2022 年末，ChatGPT 突然爆火，随之而来的是，生成式人工智能和 GPU 计算走向了主流。

利用生成式人工智能的热潮

几乎在一夜之间，对人工智能计算的需求急剧飙升，它成为了价值数十亿美元产业、消费级应用以及企业竞争战略的基础。大型科技公司和风险投资公司向人工智能研究初创企业和资本支出建设投入了数十亿美元，而这些资金最终都流向了英伟达，因为只有它能够满足不断激增的计算需求。

随着对人工智能计算需求的激增，企业面临着一个严峻的现实：训练和部署生成式人工智能模型的成本高得惊人。每一点效率的提升，无论多么微小，在大规模应用时都能转化为巨大的成本节约。由于英伟达的硬件已经在数据中心根深蒂固，人工智能公司面临着一个艰难的抉择：针对 CUDA 进行优化，否则就会落后。几乎一夜之间，整个行业都转向编写特定于 CUDA 的代码。结果是，人工智能的突破不再仅仅由模型和算法驱动，现在还取决于从经过 CUDA 优化的代码中榨取每一分效率的能力。

以 FlashAttention - 3 为例，这项前沿的优化技术大幅降低了运行 Transformer 模型的成本，但它是专门为 Hopper GPU 打造的，通过确保只有在英伟达最新的硬件上才能获得最佳性能，进一步强化了英伟达的技术锁定。持续的研究创新也遵循着同样的模式，比如 DeepSeek 直接采用 PTX 汇编语言，在尽可能低的层面上实现对硬件的完全控制。随着英伟达新的 Blackwell 架构即将推出，可以预见整个行业又要重新编写所有代码了。

强化 CUDA 主导地位的循环

这个系统正在加速发展并形成自我强化的态势。生成式人工智能已成为一股不可阻挡的力量，推动着对计算能力的无尽需求，而英伟达掌握着所有的优势。庞大的用户基础确保了大多数人工智能研究都是基于 CUDA 进行的，这反过来又促使人们对优化英伟达的平台进行投资。

英伟达每一代新硬件都带来新的功能和更高的效率，但这也需要重新编写软件、进行新的优化，并且对英伟达的技术堆栈产生更深的依赖。未来似乎不可避免：在这个世界里，CUDA 对人工智能计算的掌控只会越来越紧。

然而，CUDA 并非完美无缺。

巩固 CUDA 主导地位的这些因素，也正逐渐成为瓶颈，带来技术挑战、效率低下的问题，还阻碍了更广泛的创新。这种主导地位真的对人工智能研究领域有益吗？CUDA 是对开发者有利，还是仅仅对英伟达有利呢？

让我们退一步思考：我们已经了解了 CUDA 是什么以及它为何如此成功，但它真的好吗？我们将在下面章节探讨这个问题。

CUDA占据主导地位，但它真的好吗？

回答 CUDA 是否 “好” 这个问题，远比听起来要棘手。我们讨论的是它的原始性能？它的功能特性？还是它在人工智能开发领域更广泛的影响呢？

CUDA 是否 “好”，这取决于你问的是谁以及他们的需求是什么。在本章节中，我们将从那些日复一日使用它的人，也就是在生成式人工智能（GenAI）生态系统中工作的人的角度来评估 CUDA：

1. 对于在 CUDA 基础上进行开发的人工智能工程师来说，它是必不可少的工具，但同时也伴随着版本管理的麻烦、驱动行为不透明以及对平台的深度依赖等问题。

2. 对于为英伟达硬件编写 GPU 代码的人工智能工程师而言，CUDA 提供了强大的优化能力，但要获得顶级性能，就必须承受其中的痛苦。

3. 对于那些希望自己的人工智能工作负载能在多家厂商的 GPU 上运行的人来说，CUDA 与其说是解决方案，不如说是一个障碍。

4. 还有英伟达公司本身，它围绕 CUDA 积累了巨额财富，获取了巨额利润，并巩固了其在人工智能计算领域的主导地位。

那么，CUDA 到底 “好不好” 呢？让我们深入探讨每个角度来寻找答案！

AI工程师

如今，许多工程师在人工智能框架（如 LlamaIndex、LangChain 和 AutoGen 等智能库）的基础上构建应用程序，而无需深入了解底层硬件细节。对于这些工程师来说，CUDA 是强大的助力。它在行业内的成熟度和主导地位带来了显著优势：大多数人工智能库都设计为能与英伟达硬件无缝协作，而且大家都聚焦于单一平台，促进了整个行业的合作。

然而，CUDA 的主导地位也带来了一系列长期存在的挑战。其中最大的障碍之一是管理不同 CUDA 版本的复杂性，这简直是一场噩梦。这种无奈成为了众多梗图的主题：

这可不只是个梗图，对许多工程师来说，这是他们真实的经历。这些人工智能从业者需要不断确保 CUDA 工具包、英伟达驱动和人工智能框架之间的兼容性。不匹配的情况可能会导致令人沮丧的构建失败或运行时错误，无数开发者都有过切身体会：

“我无法使用最新的英伟达 PyTorch Docker 镜像构建系统。原因是通过 pip 安装的 PyTorch 是基于 CUDA 11.7 构建的，而容器使用的是 CUDA 12.1。” 或者：“管理英伟达 GPU 驱动和 CUDA 开发软件可能很有挑战性。升级 CUDA 版本或更新 Linux 系统可能会导致诸如 GPU 驱动损坏等问题。”

遗憾的是，这类麻烦并不罕见。解决这些问题通常需要深厚的专业知识和耗费大量时间进行故障排查。英伟达依赖不透明的工具和复杂的设置流程，这让新手望而却步，也减缓了创新的速度。

为应对这些挑战，英伟达历来都是在更高层级提供个别针对性的解决方案，而不是解决 CUDA 层本身这一根本问题。例如，它最近推出了 NIM（英伟达推理微服务），这是一套容器化的微服务，旨在简化人工智能模型的部署。虽然这可能会简化某一特定用例，但 NIM 也将底层操作抽象化了，增加了用户对其技术的依赖，限制了对底层优化和创新的接触，而这些恰恰是 CUDA 价值主张的关键所在。

在 CUDA 基础上构建应用的人工智能工程师面临着兼容性和部署方面的挑战，而那些更接近硬件底层工作的人员，即人工智能模型开发者和性能工程师，则要应对一系列截然不同的权衡取舍。

AI模型开发者和性能工程师

对于致力于突破人工智能模型极限的研究人员和工程师来说，CUDA 既是必不可少的工具，也是令人沮丧的限制因素。对他们而言，CUDA 不是一个简单的 API，而是他们编写的每一个对性能至关重要的操作的基础。这些工程师在底层进行优化工作，编写自定义 CUDA 内核，调整内存访问模式，尽可能地从英伟达硬件中榨取每一分性能。生成式人工智能的规模和成本要求他们这么做。但 CUDA 是赋予了他们能力，还是限制了他们的创新能力呢？

尽管 CUDA 占据主导地位，但它也逐渐显露出不足。它设计于 2007 年，远在深度学习出现之前，更不用说生成式人工智能了。从那以后，GPU 发生了巨大的变化，张量核心（Tensor Cores）和稀疏特性成为人工智能加速的核心要素。CUDA 早期的贡献是让 GPU 编程变得容易，但它并没有随着现代 GPU 为实现 Transformer 和生成式人工智能性能所必需的特性而发展。这就迫使工程师们不得不绕过它的局限性，才能获得工作负载所需的性能。

CUDA 无法充分发挥现代 GPU 的全部功能

像 FlashAttention-3（示例代码 ）和 DeepSeek 的创新技术等前沿技术，要求开发者绕过 CUDA，使用英伟达更低级的汇编语言 PTX。PTX 的文档并不完善，在不同硬件代际之间不断变化，对开发者来说实际上就是个黑箱。

更麻烦的是，PTX 比 CUDA 更依赖英伟达，而且可用性更差。然而，对于追求前沿性能的团队来说，别无选择，他们只能绕过 CUDA，承受诸多不便。

张量核心：提升性能的关键，但使用难度大

如今，人工智能模型的大部分浮点运算（FLOPs）来自 “张量核心”，而非传统的 CUDA 核心。然而，直接对张量核心进行编程并非易事。虽然英伟达提供了一些抽象工具（如 cuBLAS 和 CUTLASS），但要充分发挥 GPU 的性能，仍需要晦涩难懂的知识、反复的试验和测试，而且常常需要逆向工程一些未文档化的行为。随着每一代新 GPU 的推出，张量核心都会发生变化，但相关文档却未能及时更新。这使得工程师在充分挖掘硬件潜力时资源有限。

人工智能用 Python，而 CUDA 用 C++

另一个主要限制是，编写 CUDA 代码基本上需要使用 C++，而现代人工智能开发大多是用 Python 完成的。使用 PyTorch 进行人工智能模型和性能开发的工程师并不想在 Python 和 C++ 之间来回切换，这两种语言的编程思路差异很大。这种不匹配减缓了迭代速度，造成了不必要的阻碍，还迫使人工智能工程师在本应专注于模型改进的时候，却要去考虑底层性能细节。此外，CUDA 对 C++ 模板的依赖导致编译时间极长，而且经常会出现难以理解的错误信息。

如果你乐于专门为英伟达硬件进行开发，就会面临上述这些挑战。但如果你关注的不只是英伟达呢？

开发可移植软件的工程师和研究人员

并非所有人都愿意开发锁定英伟达硬件的软件，其中的挑战显而易见。CUDA 无法在其他厂商的硬件上运行（比如我们口袋里的手机芯片这类硬件 ），而且没有其他方案能在英伟达硬件上提供与 CUDA 相同的性能和功能。这就迫使开发者针对多个平台多次编写人工智能代码。

在实际操作中，许多跨平台人工智能开发工作都困难重重。早期版本的 TensorFlow 和 PyTorch 有 OpenCL 后端，但在功能和速度上都远远落后于 CUDA 后端，这使得大多数用户还是选择使用英伟达的产品。维护多条代码路径（CUDA 用于英伟达，其他方案用于其他平台）成本高昂，而且随着人工智能的快速发展，只有大型机构才有资源进行这样的工作。

CUDA 造成的这种分化形成了一个自我强化的循环：由于英伟达拥有最大的用户群体和最强大的硬件，大多数开发者首先会针对 CUDA 进行开发，并期望其他方案最终能赶上来。这进一步巩固了 CUDA 作为人工智能默认平台的主导地位。

接下来，我们探讨 OpenCL、TritonLang 和 MLIR 编译器等替代方案，并了解为什么这些选项没有对 CUDA 的主导地位造成影响。

CUDA 对英伟达自身有好处吗？

答案当然是肯定的：“CUDA 护城河” 造就了赢者通吃的局面。到 2023 年，英伟达占据了数据中心 GPU 市场约 98% 的份额，巩固了其在人工智能领域的主导地位。正如我们在之前的文章中所讨论的，CUDA 是连接英伟达过去和未来产品的桥梁，推动了像 Blackwell 这样的新架构的应用，维持了英伟达在人工智能计算领域的领先地位。

然而，像Jim Keller这样的传奇硬件专家认为，“CUDA 是片泥沼，而非护城河”，他将其与拖累英特尔的 X86 架构相类比。

Jim Keller认为， “ CUDA 是沼泽，而不是护城河”

CUDA 怎么会成为英伟达的问题呢？这存在几个挑战。

CUDA 的可用性对英伟达影响最大

黄仁勋曾说过，英伟达雇佣的软件工程师比硬件工程师还多，其中很大一部分人都在编写 CUDA 相关代码。但 CUDA 在可用性和可扩展性方面的挑战减缓了创新速度，迫使英伟达大量招聘工程师来解决这些问题。

CUDA 的庞大体量影响新硬件的推出

CUDA 在英伟达自身不同代际的硬件之间无法实现性能的可移植性，而且其庞大的库规模是一把双刃剑。当推出像 Blackwell 这样的新一代 GPU 时，英伟达面临一个选择：重写 CUDA，或者推出无法充分发挥新架构性能的硬件。这就解释了为什么每一代新硬件在推出时性能都不是最优的。扩展 CUDA 的功能既昂贵又耗时。

创新者的困境

英伟达对向后兼容性的坚持，这曾是 CUDA 早期的卖点之一，如今却成了 “技术债务”，阻碍了其自身的快速创新。虽然对老一代 GPU 的支持对开发者群体至关重要，但这迫使英伟达优先考虑稳定性，而非进行革命性的变革。这种长期支持耗费时间和资源，可能会限制其未来的灵活性。

尽管英伟达向开发者承诺会保持连续性，但 Blackwell 如果不打破与 Hopper PTX 的兼容性，就无法实现其性能目标，现在一些 Hopper PTX 操作在 Blackwell 上无法运行。这意味着那些绕过 CUDA 而选择 PTX 的高级开发者可能需要为下一代硬件重写代码。

尽管存在这些挑战，但英伟达在软件方面的出色执行和早期的战略决策，为其未来的增长奠定了良好的基础。随着生成式人工智能的兴起，以及基于 CUDA 构建的生态系统不断发展，英伟达有望继续处于人工智能计算的前沿，并已迅速成长为全球最具价值的公司之一。

CUDA 的替代方案在哪里？

总之，CUDA 既是恩赐也是负担，这取决于你处于生态系统的哪一方。它的巨大成功推动了英伟达的主导地位，但它的复杂性、技术债务以及对供应商的锁定，给开发者和人工智能计算的未来带来了重大挑战。

随着人工智能硬件的快速发展，一个自然而然的问题出现了：CUDA 的替代方案在哪里？为什么其他方案还没有解决这些问题呢？在接下来的章节中，我们将探讨最具代表性的替代方案，分析阻碍它们突破 CUDA 护城河的技术和战略问题。

像 OpenCL 这样的 CUDA C++ 替代方案怎么样？

生成式人工智能（GenAI）或许是新事物，但 GPU 可不是！多年来，许多人尝试用 C++ 创建可移植的 GPU 编程模型，从 OpenCL 到 SYCL，再到 OneAPI 等等。这些本是最有可能替代 CUDA 的方案，旨在实现人工智能计算的民主化，但你可能从未听说过它们，因为它们在人工智能领域并未发挥重要作用。

这些项目都为计算领域做出了有意义的贡献，但如果我们真想为未来解锁人工智能计算，就必须认真审视那些阻碍它们发展的错误，而不能只盯着成功之处。从宏观层面看，问题源于 “开放式竞合” 的挑战，即行业参与者既合作又竞争，以及过程中特定的管理失误。让我们深入探讨一下。

CUDA C++ 的替代方案：OpenCL、SYCL 等

有许多项目致力于实现 GPU 编程，其中我最了解的是 OpenCL。和 CUDA 一样，OpenCL 旨在为程序员提供类似 C++ 的编程体验，以便在 GPU 上运行代码。这背后还有我的一段个人经历：2008 年，我是苹果公司负责实现 OpenCL 的首席工程师之一（这是我当时构建的 Clang 编译器首次投入实际应用）。在我们完成开发后，做出了一个关键决定，将其贡献给 Khronos Group，以便在整个行业得到采用和标准化。

这一决定使得 OpenCL 在行业内得到广泛应用（见相关标识 ），尤其在移动和嵌入式设备领域。如今，它依然非常成功，为安卓等平台以及数字信号处理器（DSP）等专业应用中的 GPU 计算提供支持。与 CUDA 不同，OpenCL 从一开始就设计为具有可移植性，旨在支持 CPU、GPU 和其他加速器之间的异构计算。OpenCL 还启发了其他系统，如 SyCL、Vulkan、SPIR-V、oneAPI、WebCL 等等。

然而，尽管 OpenCL 在技术上有优势且应用广泛，但它从未成为主导的人工智能计算平台。主要原因有以下几点：开放式竞合中固有的矛盾、由此产生的技术问题、人工智能不断变化的需求，以及英伟达针对 TensorFlow 和 PyTorch 的统一战略。

委员会决策速度下的 “竞合” 问题

2008 年，苹果在个人电脑领域还只是个小角色，当时认为行业标准化能让其接触到更多开发者。虽然 OpenCL 确实在硬件制造商中得到广泛采用，但其发展很快遇到了一个重大障碍：委员会驱动的开发速度太慢。对苹果来说，这种缓慢、依赖共识的决策过程是个大问题：我们希望快速推进平台发展，添加新功能（比如添加 C++ 模板），并展现苹果平台的差异化。我们面临着一个残酷的现实 —— 委员会制定标准的弊端在于，一切都按照委员会达成共识的速度推进，而这个速度就像冰川移动一样缓慢。

硬件供应商们认识到统一软件生态系统的长期益处，但在短期内，他们又是激烈的竞争对手。这就引发了一些微妙却很严重的问题：参与者们不会向委员会透露正在研发的硬件特性（以免让竞争对手抢占先机），而是会在硬件发布后才公开这些创新，并且只有在这些特性成为通用功能后才会进行讨论（转而使用特定供应商的扩展）。

合作竞争：竞争对手之间的“合作”

这种情况对苹果来说是个大麻烦，因为苹果想要快速秘密推进项目，以便在产品发布时大放异彩。因此，苹果决定放弃 OpenCL，推出了 Metal 替代它，从未在 iOS 系统中引入 OpenCL，后来还在 macOS 系统中弃用了它。其他公司虽然继续使用 OpenCL，但这些结构性挑战持续限制着它，使其无法跟上前沿人工智能和 GPU 创新的步伐。

OpenCL 的技术问题

虽然苹果大胆地将 OpenCL 标准贡献给了 Kronos，但并没有全力以赴：苹果贡献的只是 OpenCL 的技术规范，而没有完整的参考实现。尽管编译器前端（Clang）的部分是开源的，但没有共享的 OpenCL 运行时，这就迫使供应商们开发自己的定制版本并完善编译器。每个供应商都必须维护自己的实现（即 “分支”），由于缺乏共享且不断发展的参考标准，OpenCL 变成了一个由特定供应商的分支和扩展拼凑而成的产物。这种碎片化最终削弱了它原本旨在实现的可移植性。

此外，由于供应商们隐瞒差异化特性，或者将这些特性分离成数量众多的特定供应商扩展，导致 OpenCL（及其衍生项目）碎片化，侵蚀了它作为一个统一的、与供应商无关的平台的能力。

OpenCL 在兼容性和一致性测试方面的不足，更是加剧了这些问题。最重要的是，它还存在我们之前提到的所有 “C++ 相关问题”。

开发者们想要的是稳定且得到良好支持的工具，但 OpenCL 的碎片化、薄弱的一致性测试以及不一致的供应商支持，让使用它成为了一件令人沮丧的事。一位开发者总结说，使用 OpenCL “就像拥抱仙人掌一样难受”！太扎心了。

一位开发人员将使用 OpenCL 描述为“就像拥抱仙人掌一样难受”。

就在 OpenCL 受困于碎片化和缓慢的发展速度时，人工智能在软件框架和硬件性能方面都在迅速进步。这使得 OpenCL 所能提供的功能与现代人工智能工作负载的需求之间的差距越来越大。

AI研究和AI GPU硬件不断变化的需求

TensorFlow 和 PyTorch 的推出掀起了人工智能研究的革命，不断完善的基础设施和大型企业大量资金的涌入为其提供了强大动力。这给 OpenCL 带来了巨大挑战。虽然 OpenCL 支持 GPU 计算，但缺乏大规模训练和推理所需的高级人工智能库和优化。与 CUDA 不同，它没有对矩阵乘法、Flash Attention 或数据中心规模的训练等关键操作的内置支持。

尽管将 TensorFlow 和 PyTorch 扩展以使用 OpenCL 的跨行业努力有明显的需求，但很快就遇到了根本性的障碍。那些坚持使用 OpenCL 的开发者很快发现了一个残酷的现实：如果无法充分发挥新硬件的性能，那么对新硬件的可移植性就毫无意义。由于无法实现可移植的特定硬件增强功能，再加上竞合关系破坏了合作，相关进展陷入了停滞。

一个明显的例子是，OpenCL 至今仍未对张量核心（Tensor Cores）提供标准化支持，而张量核心是现代 GPU 和人工智能加速器中实现高效矩阵乘法的专用硬件单元。这意味着，与使用 CUDA 或其他碎片化的特定供应商原生软件相比，使用 OpenCL 通常会导致性能降低 5 到 10 倍。在生成式人工智能领域，计算成本已经高得惊人，性能降低 5 到 10 倍可不只是不方便的问题，而是根本无法接受。

英伟达针对 TensorFlow 和 PyTorch 的战略举措

当 OpenCL 在碎片化管理的重压下艰难前行时，英伟达采取了截然不同的策略。正如我们前面讨论过的，英伟达的策略是严格控制、极具战略性且卓有成效的。英伟达积极与 TensorFlow 和 PyTorch 共同设计 CUDA 的高级库，确保这些框架在英伟达硬件上始终能达到最佳运行效果。由于这些框架原生就构建在 CUDA 之上，英伟达一开始就占据了巨大优势，并且通过优化使其性能从一开始就出类拔萃。

英伟达虽然保留了 OpenCL 的实现，但从战略上对其进行了限制（比如无法使用张量核心），这就确保了 CUDA 的实现始终是必要的。随着英伟达在行业内的主导地位不断巩固和提升，它不断加大对 CUDA 实现的投入。久而久之，OpenCL 的支持逐渐减少，直至消失，而 CUDA 则巩固了其作为无可争议的行业标准的地位。

我们能从这些 C++ GPU 项目中学到什么？

经历过这些的人都很清楚上述历史，但真正的价值在于从过去中吸取教训。基于此，我认为成功的系统必须具备以下几点：

• 提供参考实现，而不只是一份书面规范和 “兼容性” 测试。一个可用、可采用且可扩展的实现才应定义兼容性，而不是一份 PDF 文档。

• 由维护参考实现的团队提供强有力的领导和明确的愿景。

• 在行业领先企业的硬件上实现顶级性能，否则它永远只能是二流的替代方案，无法成为统一行业的标准。

• 快速发展以满足不断变化的需求，因为人工智能研究不会停滞不前，人工智能硬件创新仍在加速。

• 提供出色的可用性、工具和快速的编译时间，以此赢得开发者的青睐。另外，在人工智能领域，“类似 C++” 可不是什么卖点！

• 建立开放的社区，因为没有广泛的采用，技术实力再强也无济于事。

• 避免碎片化，一个分裂成不兼容分支的标准，无法为软件开发人员提供有效的统一框架。

这些就是我认为像 OpenCL 这样由委员会推动的项目永远不会成功的根本原因。这也是为什么我对英特尔的 OneAPI（现 UXL Foundation）等项目更加怀疑，这些项目名义上是开放的，但实际上由单一硬件供应商掌控，而该供应商又与其他所有供应商存在竞争关系。

AI编译器呢？

在 C++ 相关方案未能为硬件制造商统一人工智能计算的同时，人工智能行业面临着一个更大的挑战，即使是在英伟达硬件上使用 CUDA 也存在问题。如果所有代码都需要人工编写，我们要如何实现人工智能计算的扩展呢？芯片种类繁多，人工智能算法层出不穷，工作负载的组合更是不计其数，靠人工优化根本无法完成。

随着人工智能的影响力不断扩大，它不可避免地引起了系统开发者和编译器工程师（包括我在内）的关注。在下一篇文章中，我们将深入探讨广为人知的 “人工智能编译器” 栈，如 TVM、OpenXLA 和 MLIR，分析哪些成功了，哪些失败了，以及我们能从中吸取什么教训。遗憾的是，这些教训与上面提到的并没有太大不同：

历史不会简单重复，但总会押着相同的韵脚。—— 马克・吐温

人工智能编译器（TVM 和 XLA）怎么样？

在人工智能硬件发展的早期阶段，编写高性能的 GPU 代码虽然繁琐，但仍在可掌控范围内。工程师们可以用 C++ 为所需的关键操作精心编写 CUDA 内核，英伟达再将这些内核整合到像 cuDNN 这样的库中，以此巩固其行业地位。但随着深度学习的不断发展，这种方式完全行不通了。

神经网络规模越来越大，架构愈发复杂，研究人员对迭代周期的速度要求也越来越高。像 PyTorch 这样的框架中，独特算子的数量呈爆发式增长，如今已达数千个。要为每个新的硬件目标手动编写并优化每个算子？这根本不可能。

PyTorch 运算符按版本计数

这一挑战促使行业发生了根本性转变：既然手动编写内核不可行，那要是有一个编译器能自动生成内核会怎么样呢？人工智能编译器应运而生，旨在解决这一难题，这标志着从人工编写 CUDA 代码向机器生成、硬件优化计算的转变。

但历史表明，构建一个成功的编译器栈不仅是技术上的挑战，更是生态系统、碎片化和控制权的争夺战。那么，哪些成功了？哪些失败了？我们能从 TVM 和 OpenXLA 等项目中学到什么呢？

什么是 “AI编译器”？

人工智能编译器的核心是一个系统，它能够将像 PyTorch 或 TensorFlow 中的高级操作，自动转换为高效的 GPU 代码。它执行的最基本优化之一叫做 “内核融合”。为了理解其重要性，我们来看一个简单的例子：先将两个矩阵相乘（“矩阵乘法”），然后应用 ReLU（修正线性单元）激活函数。这些是常见神经网络中简单却重要的操作。

简单方法：两个独立内核

最直接（但效率低下）的做法是先进行矩阵乘法，将结果存储在内存中，然后再次读取该结果以应用 ReLU 函数。

对于可能编写 CUDA 内核的工程师来说，这些操作非常熟悉（不过要记住，CUDA 使用的是复杂的 C++ 语法！），并且有许多提高效率的实现技巧。

虽然上述方法简单且模块化，但这样执行操作的速度极慢，因为在matmul()之后，整个矩阵C被写入内存，然后在relu()中又再次读取。这种内存数据传输严重影响性能，在 GPU 上尤其如此，因为在 GPU 中，内存访问的成本比本地计算更高。

融合内核：一次遍历，无额外内存传输

解决方案很简单：我们可以将这两个操作 “融合” 为一个内核，消除冗余的内存访问。在matmul()之后，我们不再存储C，而是在同一循环中立即应用relu()：

虽然这种转换带来的性能提升因硬件和矩阵大小而异，但效果可能非常显著：有时性能可提升两倍！为什么会这样呢？通过融合操作：

• 我们消除了一次额外的内存写入 / 读取操作，减轻了内存带宽的压力。

• 我们将数据保留在寄存器或共享内存中，避免了缓慢的全局内存访问。

• 由于中间缓冲区被移除，我们减少了内存使用以及分配 / 释放内存的开销。

这只是内核融合的一个简单示例：还有许多更强大的转换方法，人工智能内核工程师一直在挑战优化的极限（了解更多 ）。随着生成式人工智能对计算需求的不断攀升，这些优化比以往任何时候都更加关键。

卓越性能，但复杂度呈指数级增长！

对于那些追求低成本和最前沿性能的人来说，实现这类优化既令人兴奋又充满乐趣，但背后隐藏着一个事实：这种方法无法扩展。

现代机器学习工具包包含数百种不同的 “操作”，如矩阵乘法、卷积、加法、减法、除法等，除了 ReLU 之外，还有数十种激活函数。每个神经网络需要以不同方式组合这些操作，这导致需要实现的组合数量呈爆炸式增长（数百种操作 × 数百种操作 = 多得数不清）。英伟达的库（如 cuDNN）提供了一个固定的选项列表供选择，但无法满足新研究的通用性需求。

此外，还有其他方面的复杂性增长：新的数值数据类型（如 “float8”）不断涌现，当然，人工智能需要支持的硬件种类也在激增。

复杂性的三个维度

早期人工智能编译器：TVM

人工智能编译器有很多，其中最早且最成功的之一是 TVM（Tensor Virtual Machine：张量虚拟机）。这个系统获取来自 TensorFlow/PyTorch 的模型，并针对不同硬件进行优化，即自动应用内核融合技术。该项目大约在 2016 年由华盛顿大学的陈天奇和路易斯・塞泽教授发起，2018 年一篇概述 TVM 架构的论文介绍了其多项创新成果和性能优势。TVM 后来开源并融入了 Apache 项目。

在其发展过程中，TVM 被众多硬件制造商采用（包括 ARM、高通、Facebook、英特尔等公司的公开贡献 ），应用于嵌入式、数字信号处理（DSP）等多个领域。TVM 的核心贡献者后来创立了 OctoAI，英伟达在 2024 年末收购了该公司，从而控制了许多 TVM 的原始开发者，这也可能影响该项目的未来发展。

TVM 是人工智能编译器行业的重要一步，但我们能从中学到什么呢？以下是我的关键收获。免责声明：虽然 TVM 使用了 LLVM，我也对其很感兴趣，但我从未直接参与其中。这是我作为局外人的观点。

无法在现代硬件上实现最佳性能

TVM 在现代人工智能硬件上难以实现最佳性能，尤其是随着 GPU 向张量核心（TensorCores）和其他专用加速技术发展。虽然 TVM 后来增加了对新硬件的支持，但往往滞后，无法充分释放硬件性能。因此，它和 OpenCL 有同样的问题：如果无法充分利用硬件，就无法实现高性能。

商业利益冲突导致碎片化

与 OpenCL 不同，TVM 不仅仅是一个规范，它有实际的实现。这使得它开箱即用的实用性更强，吸引了众多硬件供应商。但碎片化问题依然存在：供应商们对代码进行分叉，做出不兼容的更改，并且难以保持同步，这减缓了项目进展。这导致架构变更的执行出现摩擦（因为下游供应商抱怨他们的分叉版本被破坏），进而阻碍了开发。

需要敏捷应对人工智能的快速发展

最后一个挑战是，TVM 出现得比较早，但它周围的人工智能创新速度极快。由于得到谷歌、Meta 和英伟达等大型公司的支持，TensorFlow 和 PyTorch 迅速发展，性能不断提升，这也改变了 TVM 的性能对比基准。而生成式人工智能的出现更是给 TVM 带来了致命一击，改变了整个行业格局。TVM 是为 “传统人工智能（TradAI）” 设计的，处理的是相对简单的需要融合的算子，而生成式人工智能涉及与硬件深度集成的大型复杂算法，如 FlashAttention3。随着行业的发展，TVM 逐渐落后。

相对没那么重要（但仍然不可忽视）的是，TVM 还存在技术问题，比如由于过度自动调优导致编译时间极长。这些问题共同导致项目发展放缓。

如今，英伟达雇佣了许多 TVM 的原核心成员，这使得 TVM 的未来充满不确定性。与此同时，谷歌则通过 OpenXLA 践行自己的理念……

谷歌的 XLA 编译器：一个名称下的两个不同系统

与起源于学术项目的 TVM 不同，XLA 是由谷歌开发的。谷歌是最先进的人工智能公司之一，资金雄厚，在人工智能硬件领域有着重大利益。谷歌开发 XLA 是为了在其（如今已取得成功的）TPU 硬件中取代 CUDA，确保为自身人工智能工作负载实现紧密集成和最佳性能。2017 年，我加入谷歌大脑团队，帮助将 TPU（以及 XLA）从一个实验项目扩展为全球第二成功的人工智能加速器（仅次于英伟达）。

Google TPU

谷歌有数百名工程师参与 XLA 的开发（具体人数因统计方式而异），其发展迅速。谷歌增加了对 CPU 和 GPU 的支持，并最终成立了 OpenXLA 基金会。XLA 被用作多个重要硬件项目的人工智能编译器基础，包括 AWS 的 Inferentia/Trainium 等。

除了代码生成，XLA 最大的成就和贡献之一是能够处理大规模机器学习模型。在超大规模场景下，使用数千个芯片进行训练的能力至关重要。如今，最大的实用模型开始需要先进技术将其在多台机器上进行分区，XLA 开发出了简洁有效的方法来实现这一点。

既然投入如此巨大，为什么像 PyTorch 和 vLLM 这样的领先项目不在 GPU 上使用 XLA 呢？答案是，XLA 实际上是两个不同的项目，只是品牌名称合并了，其背后存在工程师激励结构问题、管理难题以及技术问题，这些使得 XLA 在实际应用中存在困难。

谷歌使用 XLA-TPU，而 OpenXLA 面向其他用户

需要理解的最重要一点是，XLA 有两种形式：1）内部闭源的 XLA-TPU 编译器，为谷歌的人工智能基础设施提供支持；2）OpenXLA，这是面向 CPU 和 GPU 的公共项目。这两个版本有部分代码（“StableHLO”）是共享的，但 XLA 的绝大部分代码（以及相应的工程工作）是针对谷歌 TPU 的，属于闭源专有代码，并不用于 CPU 或 GPU。如今，在 GPU 上使用 XLA 通常需要调用标准的 CUDA 库来提升性能。

这就导致了严重的激励结构问题 —— 谷歌的工程师可能想要构建一个出色的通用人工智能编译器，但他们的收入与 TPU 的性能表现挂钩。领导层没有太多动力为 GPU 或其他替代硬件优化 XLA，一切都以保持 TPU 的竞争力为目标。以我的经验来看，如果某项设计变更可能影响 TPU 性能，XLA 就不会优先考虑那些对其他芯片有利的改变。

结果就是，XLA 在 TPU 上表现出色，但在其他地方却不尽如人意。

OpenXLA 的管理

XLA 很早就作为开源项目发布，但明确由谷歌控制。谷歌凭借在人工智能领域的早期领先地位和 TensorFlow，使得 XLA 被行业内其他团队采用。2023 年 3 月，该项目更名为 OpenXLA，并宣布独立。

尽管进行了更名，但谷歌仍然控制着 OpenXLA（从其管理结构可以看出），而且似乎也没有持续投入：社区贡献不断减少，OpenXLA 的官方账号自 2023 年起就不再活跃。

XLA 的技术挑战

和 TVM 一样，XLA 也是围绕一组固定的预定义算子（StableHLO）设计的。这种方法在 2017 年处理像 ResNet-50 这样的传统人工智能模型时效果很好，但在应对现代生成式人工智能工作负载时却力不从心，因为现代生成式人工智能在数据类型、自定义内核和特定硬件优化方面需要更高的灵活性。如今，这是一个关键问题，现代生成式人工智能算法在数据类型上需要创新（见下图），或者像 DeepSeek 所展示的那样，在硬件层面和新型通信策略上进行创新。

vLLM 0.7 中按硬件类型支持的数据类型

因此，XLA（和 TVM 一样）也被生成式人工智能甩在了后面：如今，许多关键工作负载都是在像 Pallas 这样的实验性系统中编写的，即使在 TPU 上也绕过了 XLA 编译器。核心原因是，为了简化人工智能编译，XLA 对硬件进行了过多的抽象。这在早期人工智能模型中可行，但生成式人工智能需要对加速器进行细粒度控制，而这正是 XLA 天生不具备的能力。所以，和 TVM 一样，XLA 也逐渐被淘汰。

从 TVM 和 XLA 中吸取的教训

我为我们在 XLA-TPU 中取得的技术成就感到自豪：XLA 为许多代际研究突破提供了支持，包括 Transformer 的发明、无数的模型架构，以及在其他地方看不到的研究和产品扩展。显然，它是英伟达之外最成功的训练和推理硬件，支撑着谷歌众多领先的人工智能产品和技术。虽然我对 TVM 了解较少，但我非常尊重它对编译器研究、自动调优以及为许多早期人工智能系统提供支持所做出的贡献。

即便如此，我们还是能从这两个项目中学到很多。回顾从 OpenCL 项目中吸取的教训：

• “提供参考实现”：它们都提供了实用的实现，而不像 OpenCL 那样只是一个技术规范。

• “拥有强大的领导力和愿景”：它们都有明确的领导团队和背后的愿景。 然而，OpenXLA 的愿景与希望采用它的硬件团队并不一致。和许多谷歌的项目一样，其长期前景不明朗，依赖它存在风险。

• “在行业领先企业的硬件上实现顶级性能”：如果不调用 CUDA 库，XLA 和 TVM 都无法充分发挥英伟达 GPU 的性能，因此，如果没有类似的库可供调用，它们在其他人工智能加速器上的性能表现也存疑。 XLA 在 TPU 上确实展现了 TPU 硬件的强大能力以及比英伟达硬件更好的扩展性。

• “快速发展”：这两个项目都是为传统深度学习构建的，但生成式人工智能打破了它们的设想。向大规模模型、复杂内存层次结构和新型注意力机制的转变，需要新的硬件 - 软件协同设计水平，而这是它们无法应对的。 这最终使得那些希望在支持生成式人工智能的现代硬件上使用它们的人对其兴趣大减。

• “赢得开发者的喜爱”：XLA 的优势在于提供了一个简单清晰、易于理解的模型，这使得 JAX 框架等得以兴起。 TVM 技术很酷，但由于编译时间长且与流行的人工智能模型不兼容，使用体验不佳。

• “建立开放的社区”：TVM 建立了开放的社区，OpenXLA 也有此目标。因此，它们都从行业应用中受益。

• “避免碎片化”：这两个项目都没有做到 ——TVM 被下游广泛分叉和修改，XLA 的代码库从不接受对非 CPU/GPU 硬件的支持，所有支持的硬件都是下游自行添加的。

AI编译器技术的优缺点

像 TensorFlow 和 PyTorch 1.0 这样的第一代人工智能框架严重依赖手工编写的 CUDA 内核，无法适应快速发展的人工智能工作负载。作为第二代方法，TVM 和 XLA 通过自动编译解决了这一问题。然而，这样做的同时，它们牺牲了第一代框架的关键优势：自定义算法的可扩展性、对硬件的细粒度控制以及动态执行能力，而这些特性对于生成式人工智能至关重要。

除了从 OpenCL 项目中吸取的教训，我们还可以提出一些期望：

• 实现完全可编程性：如果对开发者隐藏芯片的强大功能，就无法实现人工智能的民主化。如果你花费 1 亿美元购买特定类型的 GPU 集群，你肯定希望充分释放芯片的全部性能，而不受限于简化的接口。

• 充分利用 AI 的复杂性：AI 编译器的主要优势在于，它允许开发者无需手动编写大量代码，即可扩展到 AI 的指数级复杂性（运算符、数据类型等）。这对于解锁下一代研究至关重要。

• 支持大规模应用：XLA 的变革性能力在于能够轻松扩展到多个加速器和节点。这项能力对于轻松支持最大规模、最具创新性的模型至关重要。这是 CUDA 从未真正突破的领域。

尽管这些 AI 编译器各有优劣，但它们都未能完全释放 GPU 性能或实现 AI 计算的大众化。相反，它们强化了各自为政：XLA 仍然以 TPU 为中心，而 TVM 则分裂成互不兼容的特定供应商分支。它们的失败，恰恰是 CUDA 替代方案本应获得成功的方式！

也许 Triton“语言”会拯救我们？

然而，就在这些编译器苦苦挣扎的同时，一种不同的方法正在形成。它并非试图取代 CUDA，而是旨在拥抱 GPU 编程，同时使其更具可编程性。

Triton 和新一波 Python eDSL 的出现，试图弥合 CUDA 的原始强大功能与 Python 的易用性之间的差距。在下一篇文章中，我们将深入探讨这些框架，看看它们的优势所在，不足之处，以及它们是否最终摆脱了过去的错误。

当然，你已经知道答案了。CUDA帝国依然占据主导地位。但为什么呢？更重要的是——我们能做些什么呢？

那些不记得过去的人注定会重蹈覆辙。——乔治·桑塔亚那

也许有一天，编译器技术能够在不剥夺我们能力的情况下减轻我们的痛苦。

Triton 和 Python eDSL 怎么样？

人工智能编译器面临着一个根本性的权衡：它们旨在通过抽象底层细节来提升易用性和可扩展性，但现代生成式人工智能工作负载需要可编程性和硬件控制权，才能实现顶级性能。CUDA C++ 能提供这种控制水平，但其使用难度大是出了名的。与此同时，人工智能开发大多在 Python 环境中进行，因此，行业自然试图将 GPU 编程与 Python 结合起来，以弥合两者之间的差距。

但这里有个问题：Python 无法在 GPU 上运行。为了填补这一空白，研究人员开发了嵌入式领域特定语言（eDSLs），这是一种基于 Python 的抽象语言，表面上看起来像 Python，但实际上在底层会编译成高效的 GPU 代码。其理念很简单：让工程师们在无需忍受 C++ 复杂性的情况下，就能获得 CUDA 的强大功能。但它真的有效吗？

在本章节中，我们将深入剖析 Python eDSLs 的工作原理、优缺点，并仔细研究 Triton（该领域最受欢迎的方法之一）以及其他一些工具。Python eDSLs 能否同时兼顾性能和易用性，还是说它们只是实现人工智能计算民主化道路上的又一次弯路？

什么是嵌入式领域特定语言（eDSL）？

当某个特定领域拥有独特的表达方式，能提高开发者的工作效率时，就会用到领域特定语言，其中最广为人知的或许就是 HTML、SQL 和正则表达式。“嵌入式领域特定语言（eDSL）” 是一种复用现有语言语法，但通过编译器技术改变代码运行方式的领域特定语言。eDSL 广泛应用于许多系统，从分布式计算（PySpark）到深度学习框架（TensorFlow、PyTorch），再到 GPU 编程（Triton）。

例如，PySpark 允许用户用 Python 表达数据转换操作，然后构建一个优化的执行计划，使其能在集群上高效运行。同样，TensorFlow 的tf.function和 PyTorch 的torch.fx会将类似 Python 的代码转换为优化的计算图。这些 eDSL 抽象掉了底层细节，让开发者无需具备分布式系统、GPU 编程或编译器设计方面的专业知识，就能轻松编写高效代码。

eDSL 是如何工作的？

eDSL 的神奇之处在于，它会在 Python 代码运行前捕获代码，并将其转换为可处理的形式。它们通常利用装饰器（Python 的一个特性）在函数运行前拦截函数。当你使用@triton.jit时，Python 会将函数交给 Triton 处理，而不是直接执行它。

下面是一个简单的 Triton 示例：

当 Triton 接收到这段代码时，它会将函数解析为抽象语法树（AST），该树表示函数的结构，包括操作和数据依赖关系。这种表示方式使 Triton 能够分析代码模式、应用优化，并生成执行相同操作的高效 GPU 代码。

通过复用 Python 现有的语法和工具，eDSL 的开发者可以专注于构建编译器逻辑，而无需设计一门全新的语言，也不用编写自己的解析器、语法和工具链。

eDSL 的优势

eDSL 为构建领域特定编译器的人带来了巨大优势：通过将语言嵌入 Python，开发者可以专注于编译器逻辑，而不必重新发明一门完整的编程语言。设计新语法、编写解析器和构建集成开发环境（IDE）工具是一项浩大的工程，而借助 Python 现有的语法和 AST 工具，eDSL 开发者可以跳过这些步骤，直接解决眼前的问题。

eDSL 的用户也能从中受益：Python eDSL 让开发者可以在熟悉的环境中工作。他们可以使用相同的 Python IDE、自动补全功能、调试工具、包管理器（如pip和conda）以及库生态系统。开发者无需学习像 CUDA C++ 这样全新的语言，只需用 Python 编写代码，eDSL 会在底层引导代码执行。

然而，这种便利性也伴随着重大的权衡，对于期望 eDSL 表现得像常规 Python 代码的开发者来说，可能会感到沮丧。

eDSL 面临的挑战

当然，天下没有免费的午餐。eDSL 存在一些权衡取舍，其中一些问题可能会让开发者深感困扰。

看起来像 Python，但并非真正的 Python

这是 eDSL 中最令人困惑的部分。虽然代码看起来像常规 Python，但它的行为在某些关键方面并不像 Python：

为什么会这样呢？因为 eDSL 并不是在执行 Python 代码，而是捕获并将函数转换为其他形式。它决定支持哪些结构，许多常见的 Python 特性（如动态列表、异常处理或递归）可能根本无法使用。这可能会导致一些在 Python 中本应正常工作的代码突然出现无声失败或难以理解的错误。

错误和工具限制

调试 eDSL 代码可能是一场噩梦。当代码出现故障时，你通常不会得到熟悉的 Python 友好错误信息。相反，你看到的是来自编译器内部深处的晦涩堆栈跟踪信息，几乎无法判断哪里出了问题。更糟糕的是，像 Python 调试器这样的标准工具通常根本无法使用，你只能依赖 eDSL 提供的调试功能（如果有的话）。此外，虽然 eDSL 存在于 Python 环境中，但它们不能直接使用 Python 库。

表达能力有限

eDSL 通过复用 Python 的语法来工作，这意味着它们无法引入可能对其领域有用的新语法。像 CUDA C++ 这样的语言可以添加自定义关键字、新结构或特定领域的优化，而 eDSL 则局限于 Python 的子语言，这限制了它的清晰表达能力。

最终，特定 eDSL 的质量决定了这些权衡带来的困扰程度。实现良好的 eDSL 可以提供流畅的体验，而设计不佳的 eDSL 则可能是一个充满挫折的 “雷区”，总是与开发者的预期相悖。那么，像 Triton 这样的 eDSL 是否能做到恰到好处呢？它与 CUDA 相比又如何呢？

Triton：OpenAI 用于 GPU 编程的 Python eDSL

Triton 最初是哈佛大学菲利普・蒂莱（Philippe Tillet）的一个研究项目，在从事 OpenCL 相关工作多年后，于 2019 年首次发表（详见我之前关于 OpenCL 的文章 ）。蒂莱加入 OpenAI 后，该项目获得了巨大的发展动力，PyTorch 2 决定采用它，更是让...