AI加速器的终局形态,很可能是AI CPU

江湖上有些争论，一旦起了头，便不容易停。

起初不过是三两句闲谈。茶馆角落里，有人说 GPU 才是正统，久经百战，千锤百炼，早已一统北境；也有人说 LPU 才是真命天子，出手不多，却招招直奔要害。更有些新门派，名字一个比一个响亮，LPU、NPU、XPU，字母飞来飞去，如群雄并起，仿佛谁家牌匾更长，谁家便更接近未来。

那么我也来搀和一下这个乱局，我的观点是：

AI 加速器的终局形态，很可能是 AI CPU

这里说的 AI CPU，并不是今天我们通常理解的“加了一点 AI 指令的普通 CPU”，而是另一种东西：它保留 CPU 式的通用可编程性、缓存体系和软件友好性，同时把芯片上绝大部分面积和功耗都交给大缓存和大规模脉动阵列。换句话说，它在“像 CPU 一样容易编程”和“像 TPU 一样高效做矩阵乘法”之间，不是二选一，而是把两者合在一起。

这是一颗传统 CPU 风格的芯片，但每个核心旁边挂着大规模脉动阵列；芯片上有大容量、软件可管理的 SRAM；各种内存之间、网络之间，以及 SRAM 与网络之间都支持高吞吐 DMA；芯片本身就是完整计算机，不需要额外 host 才能把事情跑起来；HBM 不追求特别大，而是把大模型的权重分散到多颗芯片上；KV cache 尽量往 SSD 和更低成本的层级上疏散；矩阵乘法单元的数值格式尽量收敛到真正必要的最小集合；网络优先考虑低成本、高带宽、近距离互联；软件上则提供 PTO (Parallel Tile Operation)后端，并允许用户用标准 C++ 和Python写 kernel 和自定义 op。

我会把芯片做成传统 CPU 样式的核心外加向量扩展，再给每个核心挂上脉动阵列。我会放大核内Buffer，把它做成 MB 级的软件管理缓存，并尽量让它在没有被显式占用时，仍然能部分充当自动缓存。

我会要求 DMA 成为一等公民：各种内存层级之间都能 DMA，网络与 SRAM 之间也能 DMA，而且尽量让 DMA 本身就能带Layout变换，转置，压缩、解压功能。

我不会把这颗芯片当成只能挂在别的主机后面的设备。它应当能自己跑 Linux，自己能挂起来gdb，perf等传统功能。能自己承接 TensorFlow / PyTorch 这类框架的运行时需要。必要的话，我会在同一个核上放一些SMT的标量线程，让它们处理操作系统、输入管线、集合通讯，内存搬运，网络协议、检索类任务等杂活。

HBM 我不会一味往大了堆。权重本来就应该通过多芯片并行去分摊，而不是试图把一切都塞进单颗芯片。

我会考虑给芯片配 SSD，并通过稀疏注意力或者 indexed attention，把一部分 KV cache 放到 SSD 这一层。

脉动阵列的乘法位宽，我不会轻易追求 16 bit 的全面兼容。我更偏向于更便宜的数值体系，例如 7 bit 乘法配更宽的累加，以保证精度。

结构化稀疏我会认真支持，甚至优先考虑比 2:4 更简洁的 1:2 方案。

网络方面，我会优先看 1D 或 2D torus 这种更便宜的拓扑，尽量通过短距离铜互联和可拼接主板，避免昂贵线缆。

软件生态上，我会优先把 PTO ISA后端做好，并允许用户用标准 C++ 和 intrinsic 写自定义 kernel，对接好tilelang，triton，CUDA也能兼容支持，而不是逼着用户学一门全新的专用语言。

第一代产品，我会优先打 inference，而不是上来就碰 training。

我的基本判断是：这条路线把真正昂贵却不必要的东西砍掉了，把真正值钱的东西——也就是可编程性、缓存、脉动阵列、软件栈——保留下来。如果软件水平足够高，这种产品在计算/功耗、计算/面积、tokens per dollar 上都应该非常有竞争力。

我先把最根本的一点讲清楚：AI 硬件的基础就是CUBE。

原因其实很简单。当前主流深度学习模型里，绝大部分计算量都落在矩阵乘法上。矩阵乘法这件事，脉动阵列是最有效率的硬件实现方式。于是，一个 AI 芯片到底是什么，本质上就是：一整套服务于脉动阵列的系统。

一颗 AI 芯片当然不只包含脉动阵列。它还有缓存、控制逻辑、向量单元、标量单元、内存控制器、网络接口等许多部分。但真正决定上限的，是脉动阵列。就像一辆车不只有发动机，可真正定义动力上限的，还是发动机。

所以我会把一个 AI 加速器理解成“让脉动阵列持续吃满数据、持续有活可干的一整套机制”。这里所谓芯片利用率，说到底，主要也是在看脉动阵列到底忙不忙。如果脉动阵列闲着，芯片大概率就是亏的。

行业里不同公司的名字很多：昇腾CUBE，NV Tensor Core、MXU、Matrix Core、AMX、SME、AME、NeuronCore，听上去像各种新物种，实际上看透以后，核心都还是脉动阵列。营销会用很多名字，但真正干活的是同一类东西。

如果未来 AI 完全换了范式，不再依赖大规模矩阵乘法，那这套判断当然可能失效。但至少在今天，以及可见的未来里，矩阵乘法仍然是王。既然如此，脉动阵列就仍然是王。

这件事是整个辩经系列里最重要的技术判断之一。

设想我要做一个 N×N 的矩阵乘法。如果我有一个 N×N 的脉动阵列，它每个周期能吞进去一条来自左边的向量和一条来自上方的向量，经过若干拍后，从右边吐出结果。为了让这个阵列维持 100% 利用率，我必须源源不断地给它喂数据。

现在把阵列边长从 N 变成 2N，会发生什么？

首先，阵列面积大约变成 4 倍，单拍能做的乘加数量也变成 4 倍。也就是说，从纯矩阵乘法吞吐看，它等价于 4 个原来的小阵列。

问题在于，围绕脉动阵列服务它的那些东西——标量控制、向量搬运、寄存器读写、调度、指令发射——并不会随着阵列边长翻倍而自动变成 4 倍。于是你会发现，一个更大的阵列，相当于用更少的外围控制成本，驱动了更多的矩阵乘法。

换句话说，阵列做大的收益，并不只是“算得更多”，更关键的是：外围成本被摊薄了。

如果把阵列从 4×4 拉到 128×128，这中间是 5 次边长翻倍。对 scalar 控制部分来说，收益大体上是 4 的 5 次方量级；对阵列外部的 vector 处理来说，收益大体上是 2 的 5 次方量级。哪怕这些只是粗略估计，量级也已经足够说明问题：更大的阵列，能显著降低阵列外部那一大坨不直接产生 matmul 的成本。

这也是为什么我会认为，在一个大的CUBE旁边，加一个CPU通用计算的能力不再昂贵， 目前看来所有的控制，包括乱序执行，同步的能力，都可以限制在10%以内。硬件演进走向更大的脉动阵列，不是什么偶然趋势，而是非常自然的经济性结果。

讲到这里，很自然就会冒出一个问题：既然阵列越大越划算，那是不是直接做一个铺满整片芯片的巨型阵列就好了？就像Groq等LPU一样。

答案是不行。

第一重限制，是芯片不可能只剩下阵列。你总得留面积给缓存、内存控制器、控制核、向量处理、网络接口这些东西。哪怕未来阵列和缓存占比越来越高，其它部分也不可能彻底消失。这就是我们说的图灵税。

第二重、也是更本质的限制，是模型本身的矩阵形状。

一个 N×N 的脉动阵列，最舒服的工作方式，是去做能把 N 这个维度真正填满的矩阵乘法。如果模型里的关键维度只有 64，而我的阵列是 128×128，那么这个阵列在很多场景下就只能做到 25% 利用率，哪怕软件做得更聪明，也不过提升到 50%，还是远远不到理想状态。

于是问题就来了：真正限制阵列规模的，不是我愿不愿意把阵列做大，而是客户模型里的矩阵乘法，尤其是 attention 的关键维度，到底有多大。

Transformer 里最主要有两种矩阵乘法：attention 和 Feed Forward。Decode和Prefill等矩阵乘的维度完全不同。

Feed Forward 通常很大，非常适合大阵列。

attention 则不同。attention 里真正最麻烦的，是那个 reduction 维，也就是 K 维，经常偏小。它小到什么程度呢？对很多模型来说，可能只是 16 到 128 这个量级。于是，如果你的阵列是 256×256，那么 FF 跑起来很香，attention 就会明显吃亏。

如果没有 attention 这个约束，我们今天看到的可能就不是 64×64、128×128、256×256 这种阵列，而是更夸张的 1024×1024 Tile Size。这也是为什么LPU可以单独把这种Tile Size单独扣出来，而把attention放到GPU上去执行。

阵列大小之外，影响脉动阵列成本的另一件大事，就是低精度的数值格式。

在外积和内积的争论中，我会选择内积。原因是我愿意把大带宽推给外部的DMA去做。而让CUBE专注于不同精度的乘累加。CUBE本质上是一堆点积阵列。问题不在于“做乘加”，而在于“拿什么数做乘加”。数值格式一变，乘法器、累加器、带宽、存储、功耗、面积，全都会变。

最早大家用 FP32。这当然好理解，也通用，但它贵。无论面积还是功耗，FP32 乘法器都很重。

后来进入 FP16 时代，乘法精度降到 16 bit，而累加继续保留更高精度。这一步巨大地降低了阵列成本，也同时把带宽和存储需求砍了一半。

再后来是 BF16。它保留了和 FP32 类似的动态范围，却把精度需求压下来了。对训练而言，BF16 相比 FP16 的最大好处在于不必做 loss scaling；对硬件而言，它通常也更省。

再往下走，就是 FP8、FP4、Int8 这类更低比特格式。这里有两个事实需要同时承认。

第一，位宽继续下降，阵列乘法器确实会更便宜，吞吐也会继续上升。

第二，收益不会无限线性增长，因为加法器、累加路径、格式转换、控制逻辑这些开销并不会跟着等比例缩小。

所以我的判断是：不要一味追求“支持一切格式”。支持太多格式，会把本来应该极其规整、极其高效的脉动阵列，做成一个复杂得多、贵得多的多格式怪物。

我更倾向于：

• 阵列内部的格式尽量简化，优先选最经济、覆盖面又够大的那几个；
• 向量和标量单元可以保留相对更通用的格式，例如FP16，BF16用来处理 softmax、normalization、控制流附近这些不适合直接上阵列的运算；
• 真正昂贵的，是阵列，所以格式复杂度一定要首先在阵列上克制。

如果以 inference 为主，我甚至认为“纯 Int8 芯片”是一个非常严肃的候选方向。今天不是算不动，而是很多客户还没完全建立起“8 bit 完全可以把活做好”的软件和模型习惯。

我对结构化稀疏的态度是非常积极的。普通稀疏之所以难上硬件，是因为它太不规则。脉动阵列喜欢规则数据流，不喜欢稀稀拉拉、位置随时乱跳的数据。但 A:B 结构化稀疏不一样。比如 2:4 的意思，是每连续 4 个数里，最多 2 个非零。这样的结构足够规则，可以非常自然地映射到脉动阵列上。

如果这一点在硬件里支持得好，它的收益是非常惊人的：

• 计算吞吐接近翻倍；
• 权重带宽接近减半；
• 权重存储接近减半；
• 而额外的硬件成本并不大。

在我看来，这几乎就是“白捡”的系统级收益。

相比成熟的 2:4，我其实更想认真研究 1:2 稀疏，配合 7 bit 整数乘法去做一种更极致、更简洁的数据格式。比如，把一位拿来表示两个位置中哪个非零，剩下 7 位表示非零值本身。这样就能形成一种非常规整、非常硬件友好的 Int7+1 sparse 方案。

当然，想把 1:2 稀疏做成真正通用、好用、不显著掉精度的方案，需要很强的量化与训练 recipe 支撑。这不是“硬件一宣布支持，世界自动收敛到最优”那么简单。但一旦软件 recipe 跑通，它在 inference 里的长期价值会非常高。

我的总体立场是：

• 权重稀疏，非常值得做；
• 激活稀疏，远没那么自然，尤其 attention 里不一定划算；
• 结构化稀疏不是今天所有人都用得好，不代表它不值得硬件支持，恰恰说明硬件和软件之间还存在巨大的协同空间。

前面已经隐约提到这个问题，现在把它单独拎出来讲。Transformer 里的 attention 和 FF，对阵列的理想尺寸根本不是同一个答案。

FF 喜欢大阵列，因为它的三个维度通常都够大。attention 尤其是 attention 的某些 matmul，往往更偏爱小一点的阵列，因为它的 K 维偏小，容易填不满大阵列。

这就意味着，如果一颗芯片上所有核心都挂同一种尺寸的阵列，那么它天然就是失衡的。

• 如果你选 NPUs 那样的大阵列，FF 很舒服，attention 的利用率就吃亏。
• 如果你选 GPUs 那样相对小一点的阵列，attention 友好一些，但 FF 的硬件经济性又不够极致。

所以我更喜欢的思路是：这个核同时放两种处理单元。

一类的单元阵列很大，专门优先吃 FF；另一类核心的阵列较小，专门优先吃 attention。这样做的好处是，硬件终于不再强迫 attention 和 FF 共用一个并不真正适合二者的折中尺寸。

这当然会带来新问题：软件需要理解这种异构性；芯片面积比例要如何平衡；调度器怎么把工作分给两种核心；同一二进制能否统一运行。

但这些问题并没有看上去那么可怕。真正关键的一点是：如果我们重新设计一套指令集、软件和 RTL 从一开始就把“tile shape作为参数，而不是写死的常数”作为设计原则，那么很多问题会自然变简单。软件不该硬编码向量宽度，硬件也应该参数化去开发。这样做不只是为了同代异构阵列，更是为了跨代演进、测试和 FPGA 验证。

这也是CPU里面SVE的Variable Vector Length的设计原则，那我们为什么不能设计成两维的Tile寄存器呢？

前面为了把问题讲清楚，我大多数时候默认讨论的是方阵脉动阵列，也就是 N=K=M 这一类情形。但真实世界里，矩阵乘法是 N×K 乘 K×M，三维的角色完全不同。如果继续把它们硬看成一个维度，就会错过很多重要设计空间。

N 维：真正决定低 batch 能力的维度 

先看 N，也就是左矩阵的行数。

如果只从“阵列硬件长什么样”去看，很多人会觉得 N 不是阵列的属性，因为阵列真正物理铺开的通常是 K×M，N 只是“喂多少行进去”而已。但从利用率角度看，N 恰恰非常关键。

为什么？因为我不只是要把当前这一块数据算完，我还得在算当前块的时候，把下一块Right Tile预装进来。如果要让阵列持续 100% 利用率，就要求我在处理当前 Right Tile L0B 的这段时间里，有足够多的 Left Tile L0A 行可供消费。于是，N 的最小有效值，本质上取决于阵列左右两侧的数据带宽比例。

如果左右两侧带宽完全对称，那么为了把 Right Tile 的装填开销隐藏掉，通常就需要 N≈M。

如果我愿意给 Right Tile 一侧更高的带宽，比如是 Left Tile的两倍，那么为了达到满利用率，所需的 N 就可以进一步下降。

这个结论的价值很大。因为 decode attention 往往缺的不是 K 或 M，而是 N——也就是一次真正能同时处理多少 query。decode FF 也是类似问题：如果我想把 batch 压低、把 token latency 压低，那本质上也是在压 N。于是，一颗阵列如果能在更小的 N 下仍保持高利用率，它对低 batch、低延迟就更友好。

但这件事没有免费午餐。你在阵列入口给 Right Tile 侧多堆带宽，不只是阵列口要更宽，SRAM、HBM、cache、甚至更外层的数据通路都得一起提速，否则阵列口再宽也只是空转。所以，支持低 N，本质上是在为“低延迟模式”支付整条数据路径的额外成本。

这也解释了为什么像 Groq 那样的低 batch 路线会极其昂贵：它不是简单地把 batch 降下来，而是为了让极小 N 仍然不饿死计算单元，被迫把存储和片上带宽堆到非常夸张的程度。

K 维：attention 和 FF 的真正矛盾点 

再看 K，也就是 reduction 维，这其实是前面已经反复提到的核心矛盾。attention 的 K 经常偏小，FF 的 K 经常很大。

这意味着，如果我把 K 做很大，FF 很开心，但 attention 很容易填不满；如果我把 K 做小，attention 更舒服，但 FF 的外围控制与数据搬运成本就又摊不薄。

这里有一个不那么直观的方向：K 不一定非得和其它维度一样大。

我完全可以考虑一种 K 很深、M 和向量宽度没那么大的阵列。这样做的好处，是把原本需要多个阵列加外部向量加和才能完成的 reduction，更集中地压进单个阵列内部去做，从而减少外部控制和向量处理的负担。

但坏处也同样明显：只要 K、M、向量寄存器宽度不再一致，软硬件都会复杂很多。原本“所有向量都一样宽”这个巨大的简化不见了，寄存器组织、加载排布、编译器抽象、kernel 编写都会变难。

所以我对非方阵的态度不是“全面拥抱”，而是：它是值得认真探索的设计空间，但只应该在收益非常明确的地方引入。

M 维增加，某种意义上更像“多个阵列共享同一份 Left Tile输入”。它当然也有好处，因为这样能节省一部分 Left Tile 搬运和控制开销，但整体上，它不像 K 维那么关键。

如果我要把 M 做大，往往更自然的办法其实是把 K 和向量宽度也一起做大，这样整个系统仍然保持“统一向量宽度”的干净结构。那时候，本质上你做的还是更大的方阵，只是 N 未必同步放大而已。

总的来说，我更愿意把非方阵当成一种为了低延迟或特定 workload 而做的有针对性优化，而不是默认形态。

CUBE虽然高效，但它有一个让软件非常痛苦的特点：它极度需要被持续服务。

假设阵列每拍都要吃一行 L0A、一列 L0B，还要吐出一行L0C结果。那为了不让它饿着，围绕阵列的控制核和向量单元，几乎每一拍都得准备好下一份工作。光是“把数据递进去、把结果接出来”这件事，就已经非常密集；如果中间还夹杂 softmax、归一化、后处理、地址更新、同步、网络操作，压力会更大。阵列不是算得慢，恰恰相反，它算得太快了，以至于外围必须以一种近乎苛刻的节奏去服务它。只要有哪一拍没跟上，阵列就会空等，而阵列一空等，芯片经济性就立刻掉下去。

有几种办法能缓和这个问题。

方法一：做 2D 向量寄存器，也就是Tile寄存器

如果我的向量寄存器一次能表示的不只是 1 行，而是 8 行、16 行甚至更多行，那么我就不必每拍都去重新喂一行，而是一次性交给阵列一大块工作，让它先自己消化一阵。这能明显降低对控制核“每拍跟上”的要求。这也就是我们昇腾里面说的“分形”。Fractals。

但这里也有代价。寄存器一旦做得太深，遇到 decode attention 这类本来 query 数就很少的场景，或者图像初始处理这类 K 很浅的场景，就会显得很别扭。不是所有 workload 都愿意一次凑这么多行给你。

方法二：Ping-Pong Buffer时分复用

阵列前后做硬件队列，也能缓冲一部分节奏压力。控制核可以集中把一批输入打进去，然后暂时腾出手做别的事，再回来继续服务阵列。这个办法很实用，但深队列也不是白来的，它本身要吃面积、功耗、验证复杂度。

方法三：让阵列降频

这也是老办法。阵列太快，外围跟不上，那就让阵列慢一点。这样做能减轻软硬件两边的压力。这也就是为什么我们的CUBE和其他的单元可以做到不同的频率。

但一旦数值格式已经压到 8 bit 甚至更低，单位面积内的乘法密度已经很高了，再靠降频去换外围轻松，往往意味着你要用更多阵列去填满芯片面积，总体上不一定更划算。

方法四：给阵列做专门的 DMA

既然各类内存之间本来就需要 DMA，那为什么不能进一步做一种 直接面向脉动阵列的数据 DMA？这个在昇腾里面叫做MTE，Memory Transfer Engine。

也就是说，不是先把数据搬到寄存器或某一级 SRAM，再由控制核一点点送进阵列；而是允许 DMA 直接把成千上万行、列，甚至来自网络另一端的数据流，直接灌进阵列，或者直接把阵列结果送往别处。数据源可以是 SRAM、HBM、网络，甚至另一块阵列。

这个MTE的其实可以做的更通用，我们也可以把它叫做TMA （Tile Memory Access）。它通过TLOAD，TSTORE指令，把五六维的Tensor进行合轴，转置，补padding，这些可以做到很高效很细致，比GPU里面的TMA还可以进一步灵活。

它也可以直接对集群系统的网络进行直接操作，例如 all-reduce / reduce-scatter 之类流程。

当然，DMA 不能解决所有问题。只要中间涉及复杂融合、逐元素变换、非规则索引、特殊重排，数据最终还是得经过标量/向量核来加工。但是，这个架构还是会有CPU进行兜底。

讲 AI 芯片时，大家很容易把注意力全放在 HBM 和网络带宽上，好像只要把外部喂数问题解决了，阵列自然就能吃满。

我不这么看。

当脉动阵列越做越大，vector 单元越来越强，MTE / DMA 路径也越来越激进之后，真正会开始变得尖锐的，不只是“片外带宽够不够”，而是核内大带宽到底怎么供。

换句话说，问题会越来越从“怎样把数据送到芯片上”转向“数据一旦进了核，怎样在核里面高效流动，而不是在中心寄存器文件和 crossbar 上互相打架”。

这件事现在经常被说轻了。很多设计在纸面上看都能跑通：

• 阵列要读 A、读 B、写 C；
• vector 也要读 tile、写 tile；
• TMA 还要一边做搬运、一边做 gather/scatter；
• 旁边还可能有一些 scalar / reduction / post-op 在蹭数据路径。

如果这些访问最后都回到一块巨大的中心化 tile register file 上，那么你很快就会发现，核内带宽不是“有多少端口”的问题，而是一个系统级交通问题。

只要数据在核内的主流动方式还是：

• 从一个巨大的中心 TRegFile 里读出来；
• 送到执行单元；
• 再写回这个中心 TRegFile；
• 下一拍又重新从这里读；

那么无论你端口堆到多宽，最后都很容易被下面几件事拖住：

1. 端口数和 crossbar 代价迅速爆炸；
2. 同一份数据被重复读很多次；
3. 多个单元在时间上形成中心冲突；
4. 路径太长，时序和功耗开始难看；
5. 明明片上有很多 SRAM，却还在把它当“过路站”而不是“驻留地”。

所以我对“核内大带宽”的基本判断是：

不能把它理解成‘把一个大的 RF 做得更大、更宽、更多端口’。真正正确的方向，是把核内带宽问题改写成一个分层的数据流问题。

第一层：中心寄存器文件不该是每拍必经之路

我并不反对大 TRegFile。相反，我认为 tile 作为统一物理对象，非常值得有一个中心化、可重命名、可 checkpoint、可被 OoO 前端统一管理的物理载体。

但这里有一个边界必须明确：

中心 TRegFile 该是“物理 tile 的持有者”，不该是“所有高频数据流的唯一主通道”。

也就是说，它最适合做这些事：

• rename 后的物理 tile 映射锚点；
• 生命周期管理；
• 精确异常与 flush 恢复；
• 跨单元共享时的权威副本；
• spill / fill / copy / remap 的中心节点。

它不适合做的，是每个执行单元都每拍从这里狠狠干一遍再回写一遍。

因为那样一来，TRegFile 就会从“架构上很优雅的统一抽象”，退化成“性能上最倒霉的单点瓶颈”。

第二层：一份数据，应该尽量在核内驻留，而不是重复过路

核内大带宽最重要的一条原则，其实很朴素：

能复用的，就不要重读；能驻留的，就不要重放；能广播的，就不要复制。

以 matmul / cube 类操作为例，很多时候 B operand 会被反复用很多拍；某些 FF tile 也会被多个 group / block 持续消费；某些 reduction 结果又会在接下来的 vector/post-op 里马上被使用。

如果这些东西每次都从中心 TRegFile 重新读，那你消耗的不是“存储容量”，而是最宝贵的核内主干带宽。

所以，真正高质量的设计，必须默认在执行单元附近提供更近一层的驻留结构。

比如：

• cube 侧的 operand buffer；
• accumulator buffer；
• vector 侧的 lane-local / cluster-local tile window；
• tma 侧的 stream queue / gather assembly buffer；
• 某些一发多收场景下的 multicast / broadcast holding buffer。

这些结构的价值不只是“更快”，而是它们能把高频访问从中心主干剥离出去，让真正需要跨单元、跨阶段共享的数据再回到 TRegFile。

第三层：核内大带宽，本质上要靠“多级 buffer + 局部交换”，而不是“超级大 crossbar”

很多设计走到后面，都会不自觉地滑向一个危险方向：

“既然大家都要访问中心 tile 存储，那我就把 crossbar 做大一点，把端口再加一点。”

这条路不是不能走，但我非常怀疑它是不是终局。

因为只要你真把核内数据流拉开来看，就会发现不同单元对带宽的诉求并不一样：

• cube 要的是规则、连续、吞吐型的 operand collector；
• vector 要的是中等粒度、相对灵活的 tile read/write；
• tma 要的是突发型、组包型、搬运型流量；
• scalar / reduction / control 更像低带宽但低延迟敏感的附属访问。

把这些所有模式都压到同一个中心大交换矩阵上，当然“理论上统一”，但这往往不是最便宜、最稳、最能扩展的办法。

我更相信的路线，是把核内互连分成几层：

• 执行单元内部局部通路

  ：解决每拍最密集的数据流；

• 单元邻近 buffer 间的短程交换

  ：解决 producer-consumer 紧邻复用；

• 中心 TRegFile / tile store

  ：解决统一命名、统一持有、跨域共享；

• 必要时的稀疏跨域搬运

  ：由 DMA 负责，把大块数据搬到真正需要的地方。

换句话说，核内应该更像一个小型 memory hierarchy，而不是一个所有人都同时涌向的中央广场。

第四层：提供核内大带宽，不等于所有地方都追求同一种宽度

另一个常见误区是：

既然要高带宽，那就 everywhere 都做成同样的极宽路径。

这通常不划算。

因为核内各条路径的最佳宽度，本来就不一样。

举个直白的例子：

• cube 的 B operand 也许最值得做成超宽流式喂入；
• A operand 可能更需要低抖动和驻留复用；
• accumulator 回写可能根本不该第一时间回到中心 TRegFile，而应该先进本地 accum buffer；
• vector 的某些路径更关心 bank conflict 和 lane grouping，而不是单纯追求总线有多宽；
• tma 对 tile copy/gather 的要求，可能更像“多条独立搬运流”而不是“一条巨无霸总线”。

所以，核内大带宽的正确目标，不是“所有地方都做 4 KB/cycle”，而是：

让每类数据流在它最自然的层次和路径上获得足够带宽。

这听起来像一句废话，但实际上它意味着：你不能先拍脑袋规定一个“统一超宽接口”，再要求所有单元去适配；你必须先看数据流，再决定路径和宽度。

第五层：对软件来说，核内大带宽必须是“可感知但不难用”的，但是一定要Unified统一一个Tile RegFile

我前面一直在说，硬件不应该为了省一点面积，把软件接口做得很丑。核内大带宽这个问题也完全一样。

如果我为了榨干核内带宽，最后把机器做成一种只有手写汇编天才能用好的怪物，那这不是成功，这是把复杂度硬塞给软件团队。

真正好的设计，应该让软件能够清楚地利用这些层次，但不要求软件每次都亲手调 every wire。

我更偏好的方式是：

• 编译器和 kernel 库能显式表达 tile 驻留、reuse、prefetch、multicast、DMA overlap；
• 软件能知道哪些路径是“贵的中央路径”，哪些路径是“便宜的局部路径”；
• 但硬件不要求程序员每次手动规划所有 buffer 生命期和每一拍交换细节。

也就是说，核内大带宽的供给方式，最终必须能收敛成一种软件可利用的层次化抽象。

如果做不到这一点，那么再漂亮的核内带宽数字，也很可能被算法哥抛弃。

第六层：为什么这件事会越来越重要

• 阵列越来越大，核内要做大tiling，复用更多数据
• 片上 SRAM 越来越多；
• 算子融合越来越深；
• MTE / DMA 越来越想直接对接执行路径；
• 很多系统越来越不愿意把中间结果频繁打回 HBM；
• attention、FF、MoE、post-op 的组合越来越复杂。

这意味着，未来芯片的性能瓶颈，会越来越多地出现在“核内数据怎么复用”上，而不只是“HBM 有没有再多 20%”。

AI Core 的下一阶段竞争，不只是拼片外带宽，更是拼谁能把核内大带宽做成一种分层、可复用、低冲突、软件可利用的数据流体系。

谁还在把希望寄托于一个越来越大的中心 RF 和越来越粗的 crossbar，谁就很可能会在功耗、时序、可扩展性和软件复杂度上同时吃亏。

如果不把 AI 负载分开看，很多芯片设计结论都会失真。

我会把主流 Transformer 相关负载分成三类：

• training
• prefill
• decode

它们虽然都在“跑模型”，但对硬件的要求差异非常大。

Training

training 是最难做的负载。

因为它不仅要做前向，还要做反向、更新权重、跑 optimizer、处理更多数值稳定性问题，还要给研究者保留足够高的灵活性。所以 training 芯片在硬件上更复杂，在软件上更复杂，在生态上也最难做透。

如果一家公司上来第一颗芯片就想做 training 市场，我不会说不可能，但我会说风险很高。第一代产品更现实的打法，往往还是 inference。

Prefill

prefill 是模型在“读入已有上下文”的过程。它不是在生成新 token，而是在把已有输入变成 KV cache，并完成理解。

prefill 的一个巨大好处是：我可以同时处理很多 token。

因为输入文本已经摆在那里，我不需要像 decode 那样，一个 token 算完才能知道下一个 token 是什么。于是，prefill 在 attention 里的 query 数往往很多，矩阵乘法更容易填满阵列；它在整体上也更接近 compute-bound，而不是 memory-bound。

所以，prefill 对大阵列是友好的。

Decode

decode 是最麻烦的一类 inference。

原因不在于数学量突然大很多，而在于它的形状太差。

最朴素的 decode 一次只生成一个 token。于是，在 attention 里，我只有一个 query，要拿它去看很长很长的 KV cache。这会造成两个问题：

• 阵列利用率极低，因为 N 太小；
• 内存带宽压力极大，因为要读很多 key / value，但每次只配一个 query 去用它们。

所以 decode 的本质难题，不是“算不动”，而是“喂不满阵列，却又要拉很多内存”。这也是为什么市场上 output token 通常比 input token 贵得多——不是因为理论 FLOPs 差特别大，而是因为系统效率差很多。

但 decode 也不是毫无办法。至少有三类关键优化：

1. GQA

把多个 attention heads 共享一组 keys / values，本质上提高了 key 的复用率，也提高了每次矩阵乘法中的 query 数。这样阵列利用率和带宽利用率都能明显改善。

2. Speculative decoding

如果我能先用一个便宜模型猜接下来几个 token，再让主模型一口气验证并往前推进多步，那么每次 decode 实际处理的 token 数就不再是 1，而可能是 2、4、8。这样又进一步改善了阵列利用率和带宽复用。

3. 更激进的 attention 形式

比如 sparse attention、indexed attention，以及进一步压缩 KV cache 的各种方法。这些技术一旦做对，decode 对带宽和 HBM 的需求会被大幅重写。

所以我对 decode 的看法是：它今天确实最麻烦，但“decode 天然 memory-bound、天然不适合大阵列”并不是铁律。它很大程度上取决于软件和模型有没有把该做的优化做到位。

现在行业常讨论一个问题：prefill 和 decode 应不应该拆开到不同芯片上做？

默认聚合

最自然的做法，是 prefill 和 decode 在同一类芯片上混跑。数学上二者很像，只是每次处理 token 的数量差别极大，所以从框架实现上讲，共用一套 kernel 也很自然。

问题在于，这样混着跑时，系统经常会进入一种不平衡状态：

• prefill 更偏 compute-bound；
• decode 更偏 bandwidth-bound；
• 混在一起时，很容易出现某时刻一颗芯片几乎都在做 prefill，另一颗几乎都在做 decode。

这时，两边都没有达到理想利用率。

解耦

所谓 disaggregation，就是把 prefill 和 decode 分到不同芯片，甚至不同类型芯片上。这样每一边都可以针对自己的 workload 去做专门优化。

这条路当然成立，尤其是在 decode 和 prefill 差异还很大、软件又没法很好协调时。

但它也带来额外成本：

• KV cache 需要搬家；
• 系统调度更复杂；
• 产品形态更碎片化。

受控聚合

我更喜欢的概念是 managed aggregation，也就是“受控聚合”。

与其任由 prefill 和 decode 乱混，不如主动按比例把二者拼到同一批次中，让 compute-bound 的 prefill 去填补 decode 留下的计算空隙，让 decode 去利用 prefill 暂时不用满的带宽资源。这样可以让阵列和带宽两边同时更接近饱和。

这个思路的含义非常重要：如果你确信用户会用受控聚合，那么你甚至可以少配一些带宽。

当然，这是一种更冒险的产品判断。因为不是所有客户都能把调度做这么聪明。如果他们不会，你少配的那部分带宽就会直接变成性能抱怨。

所以这再一次说明，芯片设计不是孤立做出的，它深深依赖于你对客户软件成熟度的假设。

这里说的“像 CPU”，不是说回到X86一样的通用CPU处理器；也不是说今天那些只是在普通 CPU 上加一点矩阵指令的玄铁一样的AI处理器产品，就已经是终局形态。我的意思是，终局形态会具有下面这些属性：

• 像 CPU 一样通用、可编程；
• 像 CPU 一样有清晰的缓存层次与软件友好的执行模型；
• 像 CPU 一样可以自然地运行操作系统、运行复杂运行时、承载非规则任务；
• 像 CPU 一样可以对互联，网络，拓扑，进行灵活编排，容错，和系统虚拟化。
• 但同时，它又把芯片上绝大多数面积和功耗都交给脉动阵列与大缓存，从而获得 AI 所需的极致矩阵乘法效率。

也就是说，我心里的 AI CPU，不是“CPU 加一点 AI”，而是“把 AI 芯片做成真正可编程的计算机，而不是一个需要被 host 驱动的大号外设”。

真正贵的不是“像 CPU”，而是“绕开 CPU 之后重新发明一套不好用的东西”

很多人会有一种直觉：

CPU 太复杂、太通用，所以一定不适合做极致高效 AI；

于是，AI 芯片应该尽量摆脱 CPU 风格，变成更纯粹、更专用的东西。

我恰恰觉得，这种直觉经常会把账算反。

CPU 确实复杂，但 CPU 的复杂，是一种被整个行业消化过很多年的复杂。硬件工程师知道怎么设计 CPU 风格的执行单元、缓存、异常、中断、地址翻译、内存保护；软件工程师知道怎么给这种机器写编译器、运行时、库、操作系统接口。

换句话说，CPU 不是最简单的东西，但它是最被人类整体知识体系理解透彻的复杂东西之一。

如果我偏离这个方向，去发明一套“更 AI 化”的 bespoke 执行模型，那么我并没有真正消灭复杂度，我只是把复杂度从“已经被行业掌握的复杂度”，换成了“需要我自己重新承担的软件和硬件复杂度”。

而且，这种替换通常不是等价交换。它常常会让硬件表面上省一点，但让软件、编译器、kernel、调度和整条工程链条变得难用很多。

如果你芯片上最大的、最贵的部分已经是脉动阵列和缓存，那么你在外围控制部分采用 CPU 式可编程模型，往往不是浪费，而是最划算的选择。

软件已经难得离谱了，不要再给它制造额外的异形负担

AI 芯片成败的关键，远比很多硬件团队直觉里更依赖软件。

一个 AI 芯片团队，最后真正把产品从“能跑”推到“高利用率、低成本、客户可用”，靠的不是框图，而是：

• compiler
• runtime
• tiling
• software pipelining
• op fusion
• 并行调度
• 通信原语
• profiling 与 autotuning
• 模型侧配合

这套东西本来就已经很难了。

在这个前提下，如果我再去发明一种很不自然的硬件编程模型，让软件团队去适应它，那等于是把最稀缺的工程能力花在“克服硬件反人类之处”上，而不是花在真正创造系统价值的地方。

所以我反复强调：硬件不应该为了省一点外围面积，就把软件接口做得别扭。

阵列和缓存已经占了大头，外围如果做成 CPU 风格，让软件自然、好写、好测、好调精度，通常才是系统最优。

GPU 的很多成功，并不是因为 GPU 天生最适合 AI，而是因为 Nvidia 把一套本来很别扭的体系做到了极致

这一点非常重要。

很多人看到今天 AI 世界被 GPU 主导，就自然得出结论：GPU 一定是 AI 最合理的硬件形态。

我不同意。

我承认 GPU 今天最成功，我也承认 Nvidia 在软件、硬件、生态上的执行力极强。但这不等于 GPU 这个范式本身就是终局。

恰恰相反，我会说：GPU 对 AI 的胜利，很大程度上是 Nvidia 的胜利，不是 GPU 历史包袱的胜利。

GPU 身上有很多东西，本来是图形时代留下来的：

• kernel 化的执行模型
• host-device 分裂
• 大量 thread / warp 的组织方式
• warp scheduling
• path divergence
• kernel launch / switch 的各种额外问题

这些东西之所以今天没有压垮 AI，是因为 Nvidia 花了极大力气，把它们通过 CUDA、library、compiler、调优经验和生态习惯一层层包起来了。用户今天感受到的是“它很好用”，但那并不意味着它底层就是最自然的。

如果我从零开始，只为 AI 设计一台机器，我不会先假设“我们也该有 warp、kernel、host-device split，然后再想办法把这些东西变得好用”。我会反过来问：

如果没有图形历史包袱，我还会不会主动发明出这些东西？

大多数情况下，答案是否定的。

我并不是反 GPU。我只是认为，很多 GPU 特性在 AI 场景中并不应该被默认看成“正确终局”。

Kernel 模型

GPU 世界里，程序天然被切成一个个 kernel。对于图形工作负载，这历史上有它的理由；但对于越来越复杂的 AI 系统，这种分割经常只是额外负担。

如果我要 overlap compute 和 network、让一段长逻辑跨多个阶段自然衔接、或者让系统在不同时刻做不同事情，kernel 边界经常会变成麻烦，而不是帮助。

CPU 世界里，程序就是持续运行的程序，不需要频繁地“发 kernel、切 kernel、等 kernel”。这本来就是更自然的形态。

所以我会说，kernel 模型不是 AI 的理想形态，它只是 GPU 继承下来的历史结构。

Warp scheduling

warp scheduling 的本质，是当某个 warp 因为访存或其他原因停住时，硬件赶紧去跑别的 warp，把空隙填满。

这在 GPU 上当然有用，尤其因为 CUDA 世界里天然就有很多 thread / warp。

但如果换一种抽象去看，你会发现：这本质上是在用硬件实现一种软件流水化效果。

对于 AI 这种 highly structured workload，我认为很多时候完全可以用 software pipelining 达到同样目的，而且代价更低、抽象更自然。

换句话说，warp scheduling 对 GPU 是必要的，不代表对“AI 机器”是最优的。

Path divergence

这类按 lane 子集执行的机制，在某些图形和不规则 workload 上很重要。但在 AI 里，如果模型维度设计得规整，很多这类问题本来就不该出现。

我并不是说硬件一点都不该支持 tail case，而是说：不要为了迁就用户把张量维度写成 257 这种事，而在硬件里堆非常昂贵的灵活性。

AI 世界里，规整维度本来就更快、更好测、更稳定。模型侧应该承担相当一部分“别乱写尺寸”的责任。

Software-managed cache

这一点我反而是支持的。

自动 cache 很方便，但高性能 AI kernel 最后几乎总会走向显式管理。GPU 的 shared memory / software-managed cache，在这点上其实是很对的。

我对理想 AI CPU 的期待，也不是回到“只有自动 cache”，而是：

• 默认有自动 cache；
• 但允许用户或编译器把其中一部分切成 software-managed buffer。

也就是说，好的终局不是回到“纯 CPU 老式 cache”，而是把 CPU 的自然抽象和 GPU 在局部内存控制上的长处合起来。

host-device split 本身就是一种多余税负

今天的 GPU 体系，有一个常常被视为理所当然、但我始终觉得很不自然的结构：GPU 是设备，host 是另一台机器，命令、数据、同步、运行时都要在两者之间来回折腾。

这种分裂有很多额外开销：

• 要买 host；
• host 和设备之间要传命令和数据；
• 上下文切换和 kernel 调度要跨设备完成；
• failure domain 更大；
• 某些本来可以在同一台机器内部解决的事情，被迫变成“两个节点之间通信”。

如果我的 AI 芯片自己就是一台能跑 Linux 的机器，这些很多都可以消失。

我不需要先把任务送给 host，再由 host 去喂 accelerator；

我不需要为了“GPU 不是完整计算机”而单独配置一套昂贵 CPU；

我也不需要在产品定义上接受“设备必须挂在别的设备后面”这种历史遗产。

所以，从系统经济性和产品形态看，我会非常坚定地往“芯片自己就是计算机”这条路走。这本身就是 AI CPU 方向的一部分。

如果你仔细看近些年的 GPU 演化，会发现一个很有意思的趋势：

为了应对越来越复杂的 AI workload，GPU 正在不断长出越来越多“像 CPU”的能力。

比如：

• 更强的 cache / software managed memory
• 更复杂的调度与同步
• 更像线程系统的 warp specialization
• 更紧的 CPU-GPU 结合
• 更大的统一运行时和图编排能力

换句话说，随着 workload 复杂化，GPU 并不是越来越像一个纯粹的“图形式 massively threaded device”，反而越来越在往“更可编程、更可调度、更能承载复杂系统逻辑”的方向靠。

这在我看来，不是偶然，而是方向本身在说话：

当工作负载越来越复杂时，终局不会是继续忍受图形时代遗产，而是逐渐收敛到更像 CPU 的抽象。

一旦脉动阵列足够大，外围“是不是极致专用”这件事的重要性会越来越下降

真正吃面积、吃功耗、吃成本的大头，是CUBE和HBM。

只要我阵列足够大，外围控制逻辑即便不是某种极限定制、极限抠面积的奇形怪状结构，它在整体面积中的占比也会逐渐下降。

于是，继续为了那一点外围节省，而让软件和工程体系承受大量额外复杂度，越来越不划算。

这时候，最自然的结论就是：

• 阵列和SRAM继续做大、做强；朝着3D演进。
• 外围控制、执行、访存、异常、系统接口，越来越回归 CPU 式抽象；
• 整个芯片变成一个真正可编程、可跑系统、可挂阵列的 AI CPU。

如果要让我挑一个最容易被说成口号、但又最不能停留在口号层面的词，那就是 co-design。

很多团队嘴上都说软硬协同，真正落地的时候却很容易退化成一种很差的分工方式：

• 硬件先把东西做出来；
• 软件再来“适配”；
• 适配不了，就让软件同学加班想办法；
• 最后把产品做得很难用，再把问题包装成“用户还不够专业”。

这不是 co-design，这只是先后甩锅。

在我心里，真正的 co-design 是：软件和硬件一起决定架构边界。

所有软件可见的硬件接口，硬件团队都不应该在没有软件团队深度参与的情况下自行定稿。

因为硬件团队最容易犯的错误，不是算错面积，而是误判“软件忍一忍就能过去”。现实里，很多事情软件不是“不能忍”，而是“能忍，但产品会因此永久变差”。

相反，软件团队也不应该动不动就把自己的诉求直接变成硬件要求。软件想要什么，不等于硬件就应该原样给什么。真正好的 co-design，是两边都理解代价，然后一起做出克制的取舍。

我非常看重一种状态：尊重性的讨价还价。

• 这个 feature 真的是必须的吗？
• 如果不用硬件做，软件能不能优雅解决？
• 如果软件解决，代价到底是工程量、性能损失，还是长期维护负担？
• 如果硬件支持，面积、功耗、验证和时程代价到底多大？

很多决策最后不会有一个“数学上绝对正确”的答案，而是要靠判断力。也正因为如此，越不能让任何一边单独做决定。

我很反感一种做法：单独组一个“软硬件协同部”，然后让它脱离具体产品和工程链条，在半空中讨论未来。

不是说这种团队一定没用，而是它太容易丧失地面感。

真正好的 co-design，需要知道：

• 客户到底在乎什么；
• 哪些 workload 是第一优先级；
• 哪些软件路径真的有人会去写；
• 哪些功能只是 PPT 上看着漂亮；
• 哪些成本是真金白银，哪些只是理论上的优雅。

这些知识，往往来自亲手做编译器、写 kernel、调性能、配部署、谈器件价格，而不是来自抽象讨论。

所以我更相信一种“带着具体责任做 co-design”的组织方式：参与 co-design 的人，本身也应该承担真实的硬件开发、编译器开发、运行时开发或产品落地责任。真正写过自己核代码的人，这样他在做判断时，脑子里想的不是空泛理念，而是真实后果。

为什么很多 benchmark 数字其实没有意义

如果我只能给做 AI 芯片的人一个方法论提醒，那我会说：

不要迷信数字，尤其不要迷信你自己还没想清楚语境的数字。

我不是说数字不重要。恰恰相反，我是说数字非常重要，重要到你不能只看数字本身。

在 co-design 里，最常见的认知错误之一，就是看到一个实验结果，然后直接把它当成结论。问题在于，硬件和软件是联动的，一个局部改动往往会带来一堆派生后果。你如果没有把这些派生后果一起算进去，实验数字再漂亮，也经常只是幻觉。

一个典型错误：把“关掉某优化后的世界”当成真实对照组

比如有人想量化 op fusion 的价值，于是把系统里的 fusion 关掉，然后测性能差多少。结果一看，哇，差距巨大，于是得出结论：fusion 是最关键优化。

这类结论往往不成立。

因为一旦没有 fusion，系统本身的算子组织方式、IR 设计、memory movement、kernel 划分，都会跟今天完全不一样。也就是说，你拿今天这个建立在“有 fusion”前提上的系统，粗暴关掉 fusion，再去测那个“无 fusion 世界”，得到的不是一个真实世界，只是一个被故意弄残了的世界。

另一个典型错误：把软件保持不变，当成硬件变更的真实评估

比如你把脉动阵列速度翻倍，然后发现 scalar 路径跟不上，于是说“我们必须把 scalar 也翻倍”。

这也不一定对。

因为一旦硬件真的发生这种变化，编译器和 kernel 写法很可能也会随之变化。软件会更积极去做 loop unrolling、scheduling、fusion、地址计算优化。也许最后你根本不需要把 scalar 翻倍，也许只需要小改就够了。

数字最大的陷阱：它们往往是对的，但含义是错的

这是最危险的一种情况。

很多胶片、论文、产品宣讲里的数字，并不是假的，它们甚至往往是认真跑出来的。但它们经常缺了决定意义的上下文：

• 比的是不是同一种 workload？
• 比的是不是同一种并行策略？
• 对手是不是用了默认配置而不是最佳配置？
• 报的是单芯片性能、系统性能、峰值性能，还是长期稳态性能？
• 有没有把 throttling 算进去？
• tokens per second 后面的 token 质量一样吗？

所以我的态度非常直接：

数字从不自己说话。数字只有在你已经理解语境时才有价值。

在 co-design 里尤其如此，因为任何一个改动，都不只是这个改动本身，还会触发编译器、模型、调度、缓存、网络、数值格式等一连串连锁反应。

如果一个团队表面上“极度 data-driven”，但没有人认真追问这些数字到底代表什么，那这个团队很可能只是用数字包装随机决策而已。

如果我要说 AI 软件里最重要、又最容易被低估的技术，我会选 software pipelining。

因为 AI kernel 的本质几乎总是类似这样：

• 先把 tile A、tile B 从某处搬进来；
• 然后做计算；
• 然后把结果写回去；
• 接着对下一个 tile 重复。

如果你按朴素顺序做，算的时候等内存，搬的时候计算单元空着，系统利用率会惨不忍睹。

software pipelining 的意义，就是把“搬下一块、算这一块、写上一块”重叠起来，让不同硬件单元同时干活。

从理念上讲它很简单，但手写非常容易变得恶心：

• prologue / steady state / epilogue 很烦；
• 条件判断很多；
• 寄存器和缓冲区生命周期复杂；
• 一旦 tile size、数据依赖、stage 数量稍微上来，代码就很快失控。

所以我的判断很明确：

不要靠人手把 software pipelining 直接写进每个 kernel。

更合理的路线，是做一个面向 tile 的 software pipelining library，让程序员描述：

• 输入是什么 tile；
• 输出是什么 tile；
• 每一阶段做什么；
• 坐标怎样推进；

然后由库或编译器骨架来生成流水化结构。

这样一来：

• pipelining
• tiling
• op fusion
• 通信与计算 overlap

这些本来彼此纠缠的事情，就能统一进一个更高层的抽象里。

这在我看来，是 AI 软件真正应该成熟的方向。

这就是PTO Tile Libary。https://github.com/hw-native-sys/pto-isa

如果我要给一颗新 AI 芯片定义最小可行的软件接口集合，我会优先考虑三层：

• PTO backend
• C++ kernel 接口
• LLVM backend

为什么是 PTO

因为 PTO 已经是跑在现有的昇腾产品上了。

它最值钱的地方，不只是“能编译”，而是它已经帮你把 PyTorch / TileLang/ Triton/ PyPTO 那堆海量上层表达，收敛到一个更可控的中间层。对新硬件来说，这相当于直接继承一大块产业积累。

当然，你还是得自己写 backend，这事一点不轻松。但和从零写一整套“支持主流框架”的编译与 lowering 体系相比，量级已经差很多。

为什么是 C++

因为再往下，总得给用户一个真正 general 的 kernel 编程入口。

我不想再发明一门“长得像 C++ 但其实不是 C++”的新语言。我更愿意让用户直接写标准 C++，只是在里面提供必要的 intrinsic、vector type、memory primitive、DMA primitive。

这样做的好处很直接：

• 用户已有知识能复用；
• 工具链更成熟；
• 调试和测试更自然；
• 语言本身不是门槛；
• 抽象重点被迫放在库和硬件接口上，而不是放在设计一门“看上去很炫的新语言”上。

为什么是 LLVM

因为只要你想支持 C++，本质上就需要有一条靠谱的 compiler backend 路。

LLVM 不会替你解决所有问题，但它会让你少掉大量基础设施重复劳动。对一支本来就很吃紧的 AI 芯片软件团队来说，这是极现实的选择。

LLVM才是真正的虚拟指令集，以至于下面是不是灵犀指令集都无所谓了。

组织里最危险的事：让政治替代技术判断

一旦一个项目里，技术判断开始沦为政治代理，项目通常就很危险。

因为这时大家不再关心什么是对客户、对产品、对系统最好的，而开始关心：

• 哪一派赢；
• 哪个方案表面上更好讲；
• 哪些数字更容易做 PPT；
• 哪些问题可以先藏起来以后再说。

AI 芯片项目本来就已经足够复杂，如果再叠这种组织病，最后最容易发生的，就是：人人看起来都在努力，项目却在慢慢失去真正的方向。

所以在我看来，真正能把 AI 芯片做出来的团队，不只是技术上强，还必须有一种组织气质：

• 愿意面对复杂度；
• 愿意承认不确定性；
• 不拿数字当遮羞布；
• 不拿抽象愿景替代具体工程；
• 不让软件和硬件互相当对方的附属品。

也许很多年后，人们回头看今天这代 AI 芯片，会发现我们曾经在一条并不自然的道路上走了很远。我们把图形时代的遗产、工程上的妥协、软件上的无奈，以及商业上的迁就，一层一层堆在一起，然后把这一切统称为“先进架构”。可真正不会骗人的东西，其实始终只有那几个：数据是不是走了更短的路，计算是不是落在了最值钱的地方，软件是不是终于不必再替硬件收拾残局，系统是不是以更少的复杂度交付了更多的智能。

到最后，留下来的未必是今天最喧闹、最昂贵、最像未来的东西，而很可能是那个看起来最朴素、却最符合计算本性的答案。若真如此，那么 AI 芯片的终局，或许并不是我们又发明了一种多么新奇的机器，而是我们终于有勇气承认：一台真正为智能而生的计算机，本就应该长成这个样子。