AI 工程小册

主题

AI 编译器与异构算力：CUDA 的护城河是怎么松动的 ​

一句话：模型跑多快，不只取决于芯片峰值，还取决于编译器能不能把算子映射到显存层级。TileLang 这类 tile 级 DSL 把“手写 CUDA”变成“描述计算块”，同一套代码可以生成到 NVIDIA、AMD、Ascend 等后端；它松动的是软件生态锁定，不是硬件差距。

为什么需要 ​

大模型的计算不是只有矩阵乘。注意力、MoE 路由、归一化、量化、通信重叠，每一个都要落成 GPU 上的 kernel。厂商库（cuBLAS、cuDNN）覆盖了标准形状，但新模型的结构、融合模式和低精度格式往往没有现成实现；要么手写 CUDA，要么等框架补。

手写 CUDA 的问题是成本和锁定：一套代码绑定 NVIDIA 的编程模型、工具链和调试习惯；换到 AMD 的 ROCm 或昇腾的 CANN，通常要重写。CUDA 的护城河不只是“语言”，而是语言、库、工具、人才和社区叠出来的生态。AI 编译器要解决的，就是把这层映射从“人手写”变成“编译器生成”。

核心机制 ​

一、栈的分层 ​

图编译器看的是整张计算图：把相邻算子融合、选择布局、规划显存，减少 kernel 启动和中间结果。Kernel DSL 看的是一个算子内部：怎么分块、怎么用共享内存、怎么排流水。两层解决的不是同一个问题；只做图融合，遇到新的 attention 变体还是要掉回手写 kernel。

二、为什么需要 tile 级抽象 ​

GPU 的存储层级是寄存器、共享内存、L2、HBM，容量越大越慢。矩阵乘如果每个输出都从 HBM 读输入，算术强度很低，性能被带宽卡死；把矩阵切成 tile，让一个 tile 在共享内存里被复用，算术强度才能上来。

这就是 roofline 模型说的：算术强度低于机器平衡点，是内存受限；高于平衡点，才是算力受限。编译器的工作是选 tile 大小、安排 shared memory、插入异步拷贝、做 swizzle 避免 bank conflict、把循环排成流水线。TileLang 把这件事抽象成 tile 级原语：你描述“这块怎么切、怎么搬、怎么算”，它负责生成后端代码。

TileLang 由北京大学团队在 2025 年 1 月发起，构建在 Apache TVM 之上，支持 CUDA、ROCm、Metal 等后端；2026 年又增加了昇腾后端。它的价值不在“又一个 DSL”，而在把 tile 级优化从专家手感变成可复用、可迁移的描述。

三、CUDA 护城河的三个部分 ​

  • 语言与工具链：CUDA C++、PTX/SASS、Nsight 调试与性能分析；换语言容易，换调试习惯难。
  • 库与通信：cuBLAS、cuDNN、NCCL、TensorRT；这些库经过大量形状和硬件的调优，编译器短期很难全面替代。
  • 生态与人才：教程、论文复现、招聘、社区问答；迁移成本不只是代码，还有人的经验。

高层 DSL 主要侵蚀第一层：不用手写 CUDA 也能拿到不错的性能，AI 辅助代码生成进一步降低了门槛。但第二、第三层还在；所以“护城河松动”不等于“护城河消失”。

四、异构算力的现实 ​

昇腾、AMD、Metal 都有自己的后端和库；差异在算子覆盖、数值行为、通信库和调试工具。把模型从 NVIDIA 迁到另一套硬件，通常不是“改一个编译选项”，而是“逐算子验证 + 逐层对齐 + 端到端回归”。

据 TNW 与 Gigazine 2026 年 10 月的报道，TileLang 的昇腾后端已用于 DeepSeek V3.2-Exp/V4 相关适配，配套开源了六个模块；报道还称，在 128 卡昇腾 950 超节点上，prefill 达到 410 TFlops、decode 达到 360 TFlops，分别为硬件上限的 95% 和 83%。这些是厂商与媒体口径，测试形状、精度和聚类方式没有完全公开，不能直接和 NVIDIA 的峰值数字横比。

实战场景 ​

场景：8 卡推理节点上新 attention 静默回退 eager，decode 从 1800 掉到 620 tok/s ​

现象：一台 8 卡节点上的自研 MoE（带门控的新 attention）切到图编译后端发布，单机 decode 从 1800 tok/s 掉到 620 tok/s，约三分之一（经验值，需校准）；TTFT 没变，nvidia-smi 的 GPU 利用率仍是 95%，编译日志在默认级别下没有 error、没有任何告警，业务方一度以为只是“新模型更重”。

排查路径：

排除对象判定证据结论
硬件与降频nvidia-smi -q -d PERFORMANCE 时钟稳定在 boost、无 ECC/XID；单卡 4096³ GEMM micro-benchmark 达到峰值的 78%排除：卡与时钟都正常
调度与批处理引擎日志的平均 batch 24 与旧版一致，队列深度不增，TPOT 分布整体平移而非拖尾排除：不是 batch 变小或排队变长
kernel 粒度与编译日志Nsight Systems 里每步 kernel 数从 260 涨到 1900、每层 6 个以上 8 µs 级小 kernel；日志级别调到 debug 后出现 fallback to eager: op=gate_attn，该算子每 token 每层都执行确认：热路径静默回退到 eager，融合失效
对症做法：把编译日志里的回退计数做成发布门禁，热路径一旦回退就阻断，而不是等 NCU 事后发现。
python
import re, sys
log = open("compile.log", encoding="utf8").read()
fallback = set(re.findall(r"fallback to eager: op=(\S+)", log))
hot = {"gate_attn", "moe_topk", "rmsnorm"}            # 每 token 每层都走的算子
print(f"fallback={sorted(fallback)} hot={sorted(fallback & hot)}")
sys.exit(1 if fallback & hot else 0)                  # 热路径回退即阻断发布

本章生产动作里“建 kernel 覆盖矩阵”和“保留回退路径”要同时成立：回退是给长尾算子兜底的，落在每 token 都执行的算子上，等于把三分之二的损失静默化。把 compile 日志的 fallback 变成门禁，比在 NCU 里追每步 kernel 数便宜一个数量级。

复盘要点：判据是编译器“成功”的定义为热路径零回退，而不是产物能跑；默认日志级别只能证明没有 error，证明不了没有回退。红线：回退允许清单只放冷路径算子，热路径命中即阻断发布。早发现的指标是每步 kernel 数（260 → 1900 这种跳变）和 tensor core 利用率（12%），后者在 GPU 利用率 95% 的掩护下最容易被放过。

场景：128 专家 MoE 在异构卡上整图回退，prefill 只跑到 12% ​

现象：64 卡异构集群（单卡 BF16 峰值按 320 TFlops 计，经验值，需校准）迁自研 128 专家 MoE，prefill 有效算力 38 TFlops，约为单卡 GEMM 基线（峰值的 78%，约 250 TFlops）的 15%；端到端吞吐是同规模 A800 集群的 55%，HBM 吞吐 80% 而 tensor core 只有 12%。厂商口径的 410 TFlops 是 128 卡超节点、DeepSeek V3.2 形状的 prefill 数字，与本例的模型形状和并行度都不同，不能直接对标。

排查路径：

排除对象判定证据结论
硬件与互联单卡 4096³ GEMM 达峰值 78%；跨卡 all-to-all 带宽达标、无重传排除：算力与通信都不是瓶颈
精度与数值同 dtype 同 batch 与参考实现比对，最大相对误差在 1e-2 量级内排除：慢不是低精度或数值修正造成的
算子覆盖与 roofline后端 dump 的执行序列里 grouped GEMM 位置是 gather + 逐专家小 GEMM 的参考实现，覆盖矩阵该算子停在“回退”列；HBM 80% 而 tensor core 12%确认：缺融合 kernel，gather 把算术强度拉到平衡点以下
对症做法：用 tile 级描述补一个融合的 grouped GEMM，让专家维进 kernel、A 只搬一次。
python
@tilelang.jit(target="ascend", out_idx=[2])
def moe_gemm(M, N, K, E, bm=128, bn=128, bk=64, stages=3):
    with T.Kernel(T.ceildiv(N, bn), T.ceildiv(M, bm), threads=256) as (bx, by):
        As = T.alloc_shared((bm, bk), "bf16"); Bs = T.alloc_shared((bk, bn), "bf16")
        for e in T.serial(E):                     # 专家进 kernel，不再整矩阵 gather
            T.copy(A[offs[e]:offs[e] + bm, bx * bk:(bx + 1) * bk], As)
            T.copy(B[e, bx * bk:(bx + 1) * bk, by * bn:(by + 1) * bn], Bs)
            T.gemm(As, Bs, C[by * bm:(by + 1) * bm, bx * bn:(bx + 1) * bn])

本章 roofline 的判据是算术强度：朴素实现每次乘加都从 HBM 读两个元素，算术强度 0.25 FLOP/byte，与理想分块的 682.7 差三个数量级；把矩阵切成 tile、让同一块输入在 shared memory 里复用，才可能进入 295 以上的算力受限区。回退实现正停在 0.25 那一端，所以带宽打满、算力空转，补融合 kernel 才能把 12% 拉向单卡基线的 78%。

复盘要点：判据是换卡前先拿两个数——单卡 GEMM 与 attention 的可达峰值、端到端除以可达峰值；两者差一个数量级就是算子覆盖问题，不是硬件差距。红线：覆盖矩阵里热路径算子不允许停在回退列，参考实现的性能要单独标出来当性能下限。早发现的指标是“HBM 吞吐高于 70% 而 tensor core 低于 20%”这个指纹，出现即查算子表，不必等业务方比吞吐。

场景：融合三个 kernel 只换来 15%，HBM 93% 把 2.1 倍收益吃掉 ​

现象：decode 路径把 RMSNorm + QKV + RoPE 三个 kernel 融合成一个，kernel 数从 7 降到 4，单卡 decode 从 1550 到 1780 tok/s（+15%，经验值，需校准），而立项按“省掉两次 HBM 往返、有效算力 2.1 倍”预估了 2 倍；融合后 tensor core 利用率仍是 11%，HBM 跑到 3.1 TB/s，是 H100 峰值 3.35 TB/s 的 93%。

排查路径：

排除对象判定证据结论
融合是否生效kernel 数 7 → 4，中间张量不再有独立 kernel，NCU 的 dram__bytes 下降 12%排除：融合本身生效了
算力还是带宽HBM 3.1 TB/s 是峰值的 93%，同一时段 tensor core 只有 11%确认：带宽受限，不是算力不足
tile 参数与数值口径block_K 64 → 128、num_stages 2 → 3，耗时变化在 ±2% 内；与参考实现最大相对误差 1e-2 内且 dtype 未降排除：tile 没调好与精度换来的假加速
对症做法：先用 roofline 判断收益上限，再决定继续调 tile 还是改精度与 batch。
python
peak, bw = 990e12, 3.35e12                        # H100 SXM BF16
print(f"balance={peak / bw:.0f} FLOP/byte")       # 295，本章平衡点
M, w_bytes, kv_bytes = 32, 437e6, 60e6
ai = (M * w_bytes) / (w_bytes + kv_bytes)         # 本层实际算术强度
print(f"AI={ai:.0f} -> {'带宽受限' if ai < peak / bw else '算力受限'}")   # AI=28：上限是字节降幅
print(f"KV 换 fp8 + batch 64: AI={(2 * M * w_bytes) / (w_bytes + 30e6):.0f}")  # 60，仍在带宽区

本章 roofline 只给一句判据：算术强度低于平衡点就是内存受限，编译器减少的是搬运、突破不了带宽。本例 AI 约 28，比 295 低一个数量级，所以融合的收益上限就是 HBM 字节降幅 12%，实测 15% 已经贴顶；2.1 倍来自把 FLOP 比值当成时间比值。再往前走只能改分母结构：KV 换 fp8 减字节、加 batch 提 AI，两者做完 AI 也只有 60。

复盘要点：判据是带宽受限的 kernel 不存在“融合 N 个算子约等于 N 倍”的收益，上限等于 dram__bytes 的降幅。红线：以 FLOP 倍数立 kernel 项目时，验收指标必须写成“HBM 字节降幅 + 端到端 tok/s”，不写理论加速比。早发现的方式是评审时先算 AI 与 295 的关系，或看 NCU 里 dram__bytes 与 tensor pipe 利用率的比值，写代码前就能拦下。

最小上手 ​

用 4096×4096 的 FP32 矩阵乘算两个算术强度：理想分块（每个输入只从 HBM 读一次）和朴素实现（每次乘加都读两个元素）。

python
n = 4096
flops = 2 * n ** 3
ideal_bytes = 4 * 3 * n ** 2

print(f"flops={flops / 1e9:.1f} GFLOP, ideal_bytes={ideal_bytes / 1e9:.3f} GB")
print(f"tiled AI={flops / ideal_bytes:.1f} FLOP/byte")
print(f"naive AI={2 / 8:.2f} FLOP/byte")

输出：

text
flops=137.4 GFLOP, ideal_bytes=0.201 GB
tiled AI=682.7 FLOP/byte
naive AI=0.25 FLOP/byte

H100 SXM 的 BF16 峰值约 990 TFLOPS、HBM 带宽约 3.35 TB/s，机器平衡点约 295 FLOP/byte。理想分块的 682.7 高于平衡点，是算力受限；朴素实现的 0.25 远低于平衡点，是带宽受限。两者差三个数量级，这就是 tiling 的意义：不是让芯片变快，是让同一块芯片少等内存。真实 kernel 达不到理想值，但这个对比足以判断“该优化计算还是该优化搬运”。

生产动作 ​

  • 先 profile，再换编译器：用 Nsight Systems 看时间花在 kernel、通信还是启动；用 Nsight Compute 看 occupancy、带宽、tensor core 利用率。没有 profile 就换 DSL，通常只是换一种慢法。
  • 建 kernel 覆盖矩阵：算子 × 形状 × dtype × 后端，标出“厂商库覆盖 / DSL 生成 / 手写 / 回退”。新模型上线前先跑一遍，不要等线上才发现某个形状没有 kernel。
  • 数值对齐：每个生成 kernel 和参考实现比最大误差、相对误差、NaN/Inf；低精度和融合会改变累加顺序，不能只看 loss。
  • 端到端基准：分别测 prefill 和 decode，再测整条请求；kernel 快不等于服务快，通信、批处理和 KV 管理都会吃掉收益。
  • 目标硬件 CI：在真实卡型、驱动、编译器版本上跑回归；固定版本，记录每次升级的性能和数值差异。
  • 保留回退路径：编译失败或性能回退时能切回厂商库/旧 kernel；kernel 缓存要版本化，避免“昨天能跑今天不能跑”。
  • 复现厂商数字：用自己的形状、精度和并发重跑；记录测试条件，不把峰值当持续性能。

边界 ​

  • 编译器不能凭空造算力。硬件峰值、显存带宽、互联带宽是上限；编译器只能减少浪费，不能突破物理约束。
  • 峰值不等于持续性能。散热、功耗、时钟、通信都会让长时间运行低于峰值；厂商标称的 TFLOPS 通常是理想条件。
  • 可移植性有代价。一套 DSL 生成多后端，性能通常不如各后端手写极限优化；先用它覆盖长尾算子，再把最热的 20% 手写。
  • 后端成熟度不同。CUDA 的库和调试工具最完整；ROCm、CANN 在算子覆盖、数值行为和社区资料上仍有差距。
  • 生成代码难调试。性能问题可能来自 tile 大小、layout、流水，而不是模型；要保留中间代码和 profile 能力。
  • CUDA 护城河没有消失。语言层被削弱，但库、通信、工具和人才还在；短期更现实的是“多后端可用”，不是“完全替代”。

讲清楚 ​

AI 编译器把模型计算映射到硬件：图编译器做融合和布局，kernel DSL 做 tile 分块、共享内存和流水。它决定同一块芯片能跑出多少有效算力。

TileLang 这类 tile 级 DSL 把手写 CUDA 的经验变成可复用的描述，能生成到 NVIDIA、AMD、Ascend 等后端。它降低的是迁移和开发成本，不是硬件差距；最热的算子仍然需要针对后端做极限优化。

判断要不要用编译器，先算 roofline：算术强度低于机器平衡点，优化搬运；高于平衡点，优化计算。再用自己的形状做端到端基准，保留厂商库回退路径。CUDA 的护城河在松动，但还没有塌。

参考：TileLang；Apache TVM；Triton；TNW 2026-10 报道；Gigazine 2026-10 报道；NVIDIA H100

适用：多硬件后端、算子融合与推理/训练性能优化 | 初稿 2026-10

把 AI 工程讲到既懂也会用

AI 工程小册 · 持续更新