ARTICLE · INTELLIGENCE

战地情报 · 详情页

来自尧图项目组的一线实战观察与深度解析

大模型算子框架选型指南:从CUDA到Triton三层结构与实践方法

大模型算子框架选型指南:从CUDA到Triton三层结构与实践方法

搞大模型落地的人,迟早都得面对一个绕不开的词:算子。模型在 GPU 上跑得慢,你打开 profile 一看,倒在前三名的大概率就是那几个大头算子;部署时换了推理框架,速度天上地下,差别也几乎都发生在算子层。所以当同事问我“大模型算子框架到底怎么选型”,我知道他真正想解决的是:在给定硬件上,用哪套框架跑现有模型最快、最稳,同时未来加新模型时不至于被某几个算子卡死。这篇就把我的实践思路完整写出来,包括概念分层、核心算子、基准测试方法、常见坑,争取给你一张可以直接照着做的选型地图。

1. 先搞清楚你选的到底是哪一层“框架”

1.1 大模型算子框架的全貌:三层结构

我把“大模型算子框架”拆成三层来理解,这是选型的第一件事,也是最容易搞混的事。

第一层是算子运行层,也就是和硬件直接打交道的体系。NVIDIA GPU 对应 CUDA,AMD 对应 ROCm,Intel GPU 对应 oneAPI,昇腾 NPU 对应 CANN。这一层决定了一个算子能不能跑、能跑多快。比如同样的矩阵乘算子,在 CUDA 上可以调用 cuBLAS 里经过十年优化的 GEMM,在昇腾上就要走 CANN 的算子实现,两者的性能特性完全不同。

第二层是模型训练/推理框架层。训练侧常用 PyTorch、MindSpore、PaddlePaddle,推理侧常用 vLLM、TensorRT-LLM、SGLang 等。这一层负责把模型结构描述出来,然后把算子调度到第一层去执行。PyTorch 是默认选择,但推理场景下 vLLM 这类服务框架会对算子做更激进的重写和融合,所以要分开看。

第三层是算子编写与优化框架层。当你需要写一个自定义算子时,可以用 CUDA C++ 手写 kernel,也可以用 Triton 用 Python 编写,或者在 CANN 上用 Tiling 思路实现。这一层决定了你自己写算子的成本和性能上限。

选型从来不是单选某一层,而是三层组合起来选。我在实际项目里见过太多人只盯着 PyTorch 选型,最后在推理侧被 vLLM 的性能差距打得措手不及;也见过有人只谈 CUDA 算子优化,却忽略了模型框架层是否支持图融合和自动混合精度。三层一起看,才能叫完整的算子框架选型。

1.2 别把 Agent 框架、Web 框架和算子框架混为一谈

近两年“框架”这个词被用得太泛滥了,我在热词里看到 springboot、nuxt、若依、pytest、langchain、dify、crewai 都叫“框架”,但它们解决的问题完全不同。

springboot 是 Java Web 框架,nuxt 是前端框架,若依是管理后台脚手架,pytest 是测试框架,langchain、dify、crewai 是 Agent 编排框架。这些框架关心的是“功能怎么组织、服务怎么调”,不关心算子在硬件上怎么执行。一个模型服务就算用 springboot 把 HTTP 接口包得再好,底层矩阵乘能不能跑满硬件,仍然由算子框架决定。

所以我建议选型时先把边界画清楚。如果你的痛点是“模型写好了但推理延迟太高、吞吐上不去”,那你选的是算子框架层面的东西;如果你的痛点是“Agent 工作流编排复杂、工具调用链路混乱”,那 langchain 和 dify 才参与决策。两件事混在一起聊,选型会非常痛苦。

另外还要提一嘴传统视觉算子。像 HALCON 里的那些算子、拉普拉斯算子、各种滤波核,属于传统图像处理算法库范畴。虽然它们也是“算子”,但和大模型里的矩阵乘、注意力算子不是一个体系。工业视觉场景如果要用 AI 做检测,通常是大模型负责语义理解,传统算子负责几何测量,两者共存但不互相替代。

2. 选型前的三张底牌:硬件、场景和性能基线

2.1 算子性能的第一约束是硬件体系

有一个现实你必须接受:算子框架选型,本质上是硬件选型之后的第二决策。你手里是什么卡,基本就决定了你能选什么算子框架。

在 NVIDIA GPU 上,CUDA 生态成熟度无出其右。PyTorch 的算子默认走 CUDA,cuBLAS、cuDNN、FlashAttention 这些高性能算子库都是为它优化的。你在 NVIDIA 卡上选 PyTorch+CUDA,是风险最低、收益最明确的组合。

在 AMD 和 Intel 硬件上,ROCm 和 oneAPI 也在快速补齐,但算子覆盖度和性能优化深度仍然不如 CUDA。我在 AMD 卡上实测过几个主流模型,部分算子性能差距能达到 30% 以上,特别是注意力类算子和混合精度算子。这不是说不行,而是你要花更多时间去验证。

在国产 NPU 上,CANN 这类异构计算平台通常绑定自己的模型运行框架或适配层。选型时最要注意的是算子覆盖清单:你的模型里有没有 CANN 没实现的算子?如果没有,要么手写算子,要么改模型结构,这两件事成本都不低。所以在这个场景下,优先选模型结构常见的模型,比如标准的 Transformer 结构,避开冷门算子。

选型前先摸清硬件:

  • 硬件型号与算力规格,特别是 FP16/BF16/INT8 的峰值算力和显存带宽。
  • 驱动与算子库版本,CUDA 和 cuDNN 版本直接影响算子可用性。
  • 多卡互联方式,NVLink 或 PCIe,影响通信算子的选型。

2.2 训练侧和推理侧的选型逻辑完全不同

同一个团队,训练和推理选出来的算子框架组合往往很不一样。这不是拍脑袋,而是两者的优化目标差异太大。

训练侧的核心目标是吞吐量和可扩展性。你要处理的是大量样本的 forward+backward,需要自动微分、混合精度、梯度累积、ZeRO/FSDP 分布策略、张量并行这些能力。PyTorch 配合 DeepSpeed 或 FSDP,是目前最主流的组合。算子层面关心的是矩阵乘能不能吃到 Tensor Core,反向传播的 GEMM 是否被融合,通信算子能不能跟计算重叠。

推理侧的核心目标是延迟和吞吐的平衡。你要处理的是长上下文下的自回归生成,KV cache 不断增长,显存带宽成为主要瓶颈。vLLM 的 PagedAttention、continuous batching、算子融合就是针对这个场景做的。TensorRT-LLM 更是把很多算子静态编译成最优 kernel,性能很高但灵活性差一些。

我在实际项目里吃过一个亏:在训练侧用 PyTorch 默认 eager 模式跑得好好的模型,直接搬到推理服务里跑,吞吐只有理论值的四分之一。后来切到 vLLM,同样模型、同样卡,吞吐翻了两倍。原因就是 vLLM 把注意力算子重写成了分页形式,并且减少了中间张量的搬运。这就是训练推理选型逻辑不同的典型例子。

简单给一张选型对照表:

场景常见框架组合核心优化目标主要算子瓶颈
预训练/全量微调PyTorch + FSDP/DeepSpeed + CUDA吞吐量、可扩展性GEMM、AllReduce 通信
推理服务部署vLLM / TensorRT-LLM + CUDA延迟、吞吐、显存Attention、KV cache、GEMV
边缘端部署PyTorch Mobile / ONNX Runtime显存占用、功耗INT8 量化算子、轻量化算子
NPU 私有化部署对应异构框架(如 CANN)+ 模型框架算子覆盖度、国产化适配自定义算子、混合精度算子

2.3 用 roofline 和 MFU 判断算子的真实瓶颈

选型时最怕的是凭感觉说“快”或“慢”。我习惯先建立两个性能基线指标,一个是 roofline 模型,一个是 MFU。

roofline 模型解决的是“这个算子是计算密集还是访存密集”的问题。计算密集型的算子比如大矩阵乘,理论峰值算力高,主要受 FLOPS 限制;访存密集型的算子比如 LayerNorm、Softmax、Embedding,主要受显存带宽限制。判断方法是算一个算子的算术强度,也就是 FLOPs 除以 Bytes。如果这个值远高于硬件的拐点,它是计算密集;远低于拐点,它是访存密集。

举个例子,一个 4096 维的 LayerNorm,对每个 token 都要做一次均值方差归一化,计算量很小但要从显存读一遍整个向量,明显是访存密集。这就是为什么很多人把 LayerNorm 和元素级算子融合,合并到注意力或 MLP 块里,少读写几次显存,收益立竿见影。

MFU 是 Model FLOPs Utilization,衡量的是实际计算效率。测法是记录模型跑一遍的时间,估算模型的理论 FLOPs,然后除以(时间乘硬件峰值算力)。比如一块 FP16 峰值 312 TFLOPS 的卡跑 7B 模型 batch size 为 512,每次 forward 理论 FLOPs 大约 7.3e12,实测时间 0.1 秒,MFU 就是 7.3e12 / (312e12 * 0.1) ≈ 23%。这个数字低于 30% 说明算子效率还有很大提升空间,优先检查访存密集算子和 kernel launch 开销。

选型时我会把关键算子的 MFU 和带宽利用率都记下来,作为框架 A 和框架 B 的横向对比依据。数据说话,比任何宣传材料都靠谱。

3. 大模型核心算子拆解与适配要点

3.1 一眼看清单个大模型里有哪些关键算子

要选型,你得先知道模型里到底有哪些算子。以标准 Transformer 结构为例,我列一下最常遇到的算子:

  • 矩阵乘类(GEMM):QKV 投影、注意力输出投影、MLP 的 up/gate/down 投影。这是计算大头,占整个模型 FLOPs 的 60% 以上。
  • Attention 类算子:Softmax(QK^T)V、FlashAttention 的融合版本。长序列下是访存和计算的双重压力点。
  • 归一化类算子:LayerNorm、RMSNorm,访存密集,注意精度。
  • 激活函数类算子:GELU、ReLU、SiLU、SwiGLU 的变体。现在主流模型用 SwiGLU 的特别多,对算子库支持度要求高。
  • 位置编码类算子:RoPE 旋转位置编码。这东西写起来简单,但很多老的算子库没有实现,容易变成坑。
  • 嵌入类算子:Embedding、Unembedding,访存密集,词表很大时尤其明显。
  • KV cache 相关算子:存、读、分页管理,推理框架差异主要在这。

我列一个更直观的表格:

算子计算类型性能瓶颈常见优化手段
GEMM计算密集峰值算力、Tensor Core 利用率分块、算子融合、cuBLAS/tcgen05
FlashAttention计算+访存中间矩阵显存、带宽分块计算、重计算、IO 感知
LayerNorm/RMSNorm访存密集显存带宽与前后算子融合
RoPE访存密集带宽、小算子启动融合进 attention kernel
SwiGLU访存密集带宽融合在 MLP 算子中
KV cache 读写访存密集长上下文带宽PagedAttention、量化缓存

3.2 算子融合、重计算与量化:框架的隐藏竞争力

很多人选框架时只对比算子数量,但真正拉开差距的是框架对算子的“再加工”能力。同样的 TRANSFORMER 结构,不同框架跑出来的性能差一倍,秘密几乎全在下面这三件事上。

第一是算子融合。把多个小算子的中间结果留在寄存器或共享内存里,避免每次写完显存再读出来。比如 QKV 投影之后的 reshape 和 transpose,很多框架会融合进 GEMM kernel;LayerNorm 与后面的线性层融合;SwiGLU 的激活与矩阵乘融合。PyTorch 的 torch.compile 和 CUDA Graph 都在做这件事,vLLM 和 TensorRT-LLM 更是把融合玩到了极致。

第二是重计算。在训练反向传播时,为了省显存,不保存中间激活值,反向时重新算一遍。FlashAttention 就是这么干的,它不把完整的 attention 矩阵写回显存,而是在反向时重算。FlashAttention 2 和 3 还进一步优化了分块大小和线程调度。这个思路也直接影响推理框架的 KV cache 管理。

第三是量化。FP16 到 INT8、INT4 的量化,算子实现完全不一样。INT8 GEMM 需要专门的量化 kernel,保证反量化时机正确,精度不掉。AWQ、GPTQ 这些量化方法也依赖框架提供对应的量化算子。选型时一定要确认你需要的量化模式目标框架是否原生支持,而不是靠自己改算子硬凑。

这三个能力,决定了你的模型在硬件上离“理论性能天花板”有多近。我的一句话经验是:不要看框架宣传的支持算子数量,要看它对主流大模型结构的融合深度和量化支持度。

3.3 自定义算子与开发框架的选择(CUDA / Triton / CANN)

跑标准 Transformer 模型,现有算子库基本够用。但模型一改,加了新模块、新激活、新注意力变体,就绕不开写自定义算子这一步。这时候选哪个开发框架,直接影响开发效率和性能上限。

CUDA C++ 是性能上限最高的选择,也是所有其他方案的基准。你可以精确控制寄存器分配、shared memory、线程调度,能拿到理论上限。代价是开发周期长,调试痛苦,还要自己处理复杂的边界条件。对生产级算子,比如新的注意力实现,值得上 CUDA。

Triton 是更现代的选择。你用 Python 写算子逻辑,Triton 帮你做分块、调度、内存布局优化,生成的 kernel 性能大概能达到手写 CUDA 的 80% 到 95%。我写一个 elementwise 加法算子的例子:

import triton import triton.language as tl @triton.jit def add_kernel(x_ptr, y_ptr, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(axis=0) block_start = pid * BLOCK_SIZE offsets = block_start + tl.arange(0, BLOCK_SIZE) mask = offsets < n_elements x = tl.load(x_ptr + offsets, mask=mask) y = tl.load(y_ptr + offsets, mask=mask) tl.store(out_ptr + offsets, x + y, mask=mask)

这段代码性能通常接近手写 CUDA,但开发时间缩短一个量级。对大多数非极致性能下的自定义算子,我首选 Triton。

CANN 这类异构计算平台的自定义算子,通常需要你理解硬件的 Tiling 结构,开发门槛比 CUDA 还高一点。选型时我要么尽量避开非要写自定义算子的场景,要么提前预留两周算子开发时间。不要低估这个时间成本,很多项目延期就延在这。

我的完整建议是:优先用现有算子库跑通需求;新算子先用 Triton 快速实现和验证;上生产再考虑用 CUDA 重写热点算子;在 NPU 等异构平台上则直接按官方算子框架规范开发,并提前确认算子兼容清单。

4. 实操流程:一次规范的大模型算子框架选型怎么做

4.1 需求盘点与候选清单

选型不做需求盘点,就是耍流氓。我会先拉一个信息表,把技术约束全部列出来:

  • 模型结构:具体是哪几个模型,是否全部是标准 Transformer,有没有特殊算子。
  • 模型规模:参数量、上下文长度、词表大小。
  • 硬件环境:具体卡型、显存资源、是否多卡、是否评估板级边缘设备。
  • 业务指标:在线延迟要求、吞吐要求、并发数、是否需私有化部署。
  • 精度要求:FP16、INT8 还是混合,能不能接受精度损失。
  • 团队能力:团队会不会写 Triton、CUDA,还是只会用 PyTorch 训练。

候选清单也按照三层去列。运行层锁定 CUDA / ROCm / oneAPI / CANN 等;模型框架层训练侧列 PyTorch、MindSpore、PaddlePaddle,推理侧列 vLLM、TensorRT-LLM、SGLang;算子开发层列 CUDA、Triton、对应异构框架。

我会把这个清单控制在 3 到 5 个组合以内。超过 5 个组合的对比成本太高,反而不容易得出结论。比如一个典型组合是 PyTorch + CUDA + vLLM,另一个是 PyTorch + CANN + 专用推理引擎,最多再加一个 PyTorch + Triton 的自研方案。

4.2 微基准与整模型基准怎么设计

基准测试分两层:微基准测单个算子的性能,整模型基准测端到端效果。两层都要做,只做整模型无法定位瓶颈,只做微基准又脱离真实负载。

微基准设计时,我会挑选模型里耗时占比最高的 3 到 5 个算子,每个算子覆盖实际会用到的数据形状。比如 GEMM 要测 batch size 为 1 的 prompt 处理、batch size 为 64 的训练矩阵乘;Attention 要测上下文长度从 512 到 8192 的变化。每个算子至少跑 100 次,取稳定值,记录前 10 次作为预热丢弃。

整模型基准设计要贴近业务。如果是推理场景,我常用固定 prompt 长度 512,续写长度 128,测 TTFT(首 token 延迟)、TPOT(每 token 生成延迟)、吞吐和显存峰值。如果是训练场景,固定一个代表性的数据集切块,测每秒处理的 token 数和显存占用。

这里有个实操细节:跑整模型基准前,先确认框架有没有开启 CUDA Graph、continuous batching、算子融合等特性。默认配置和开启优化后的配置,性能可能差两倍。我之前对比框架时就遇到过这个坑,差点误判一个框架的性能。

4.3 性能数据怎么判读,指标怎么记

数据跑出来怎么判读,是选型中最需要经验的一环。我习惯按这三步来。

第一步看总指标。TTFT 是否在业务容忍范围,TPOT 是否流畅,吞吐是否达标,显存是否在预算内。总指标不合格,直接淘汰。

第二步看 MFU 和带宽利用率。如果某个组合的 MFU 明显低于另一个组合且场景是计算密集主导,那么基本是这个框架的算子实现不行。这时候打开对应 profiler,比如 NVIDIA 就用 Nsight Compute,逐算子看 SM 占用率、访存命中率、kernel 启动频率。

第三步看短板是否可接受。每个框架都有自己的弱项,关键是弱项是不是你的核心场景。比如 vLLM 的 PagedAttention 在很长序列下表现极好,但如果你业务全是短请求,这个优势就不重要。

指标记录我建议用一张表,字段包括算子名、数据形状、耗时、带宽利用率、显存值、是否开启融合。所有测试保持同样的随机种子和输入数据,避免因数据差异影响结果。

4.4 决策矩阵与最终结论

做完基准测试,我会把所有候选组合放进一个决策矩阵里打分。权重按业务需求定,我常用的一组是:性能 35 分、生态成熟度 25 分、团队可维护性 20 分、部署成本 10 分、灵活性 10 分。

假设有两个组合,组合 A 是 PyTorch + vLLM + CUDA,组合 B 是 MindSpore + 自研推理引擎。如果业务跑在 NVIDIA 卡上,组合 A 的性能和生态得分通常更高;如果对私有化部署和国产化环境有要求,组合 B 的部署成本和生态得分会反转。打分不是拍脑袋,每一项都要对应前面的基准测试记录。

最终结论最好写成一页纸的决策报告,内容包括:选中的组合、关键性能数据、被淘汰方案的理由、风险点、以及后续如果模型结构变化要怎么重新评估。这份报告也是你向团队其他成员解释选型原因的依据。

我在做选型报告时,还会单独写一页“如果一个月后模型变了怎么办”。模型加了视觉编码器、加了多模态输入、换了 MoE 结构,当前选型是否还成立。没有这个预案,选型很容易变成一次性决策,后面被动。

5. 常见问题与排查技巧实录

5.1 模型跑不起来的第一现场

选型落地时最经典的问题是:模型结构没问题,但新框架跑不起来,报错“算子不存在”或“not supported on this device”。我在热词里看到大量关于“大量使用算子对硬件性能的挑战”的讨论,其实算子不支持比性能慢更让人头疼。

排查思路我总结成以下步骤:

第一步,定位无法编译或无法执行的算子。PyTorch 下可以用 torch.profiler 或直接看报错堆栈,vLLM 下会打印具体算子名。Triton 或 CUDA 的报错信息越早定位越省时间。

第二步,判断该算子是否可以替换成等价算子。比如某些框架不支持 F.scaled_dot_product_attention,就把 attention 写成手写 softmax(QK^T)V 的形式,先把模型跑通,再考虑性能。

第三步,检查算子库版本。同一算子在不同版本里支持度差别很大。我遇到过 FlashAttention 算子,老版本不支持 GQA,升级后就好了。

第四步,如果必须保留该算子,就要自定义实现。评估开发成本时,先把性能要求定出来,再决定用 Triton 还是 CUDA。

一个很典型的坑是 RoPE。很多模型位置编码用的是 RoPE,但有些部署框架的算子实现里写的是绝对位置编码逻辑,导致模型精度出问题或者干脆报错。这种情况我建议优先用新版本算子库,别自己硬改。

5.2 性能差距到底来自哪里

同样硬件上,框架 A 跑模型 20 token/s,框架 B 跑 40 token/s,差距很大,但瓶颈在哪经常要找一阵子。我总结出最容易出现的三个性能坑。

第一个坑是算子没有开启融合。PyTorch eager 模式每个算子单独 launch,kernel 启动开销在短序列下尤其明显。开启 torch.compile 或 CUDA Graph 后,延迟能下降不少。推理场景下,vLLM 默认做融合,差异会体现在这里。

第二个坑是数据精度和布局不合适。默认 FP32 和 FP16 性能差一倍不止。矩阵乘的输入布局如果不符合 Tensor Core 要求,甚至会回退到普通 kernel。检查时用 profiler 看 kernel 名称,如果出现某些回退版本的名字,优先处理精度和布局。

第三个坑是显存带宽打满。长序列推理时,KV cache 读写把带宽占满,计算单元反而空闲。这时换框架也没用,得从量化 KV cache、分页管理、减少重复读取这几条路走。我在压测 32K 上下文模型时深有体会,带宽才是长文本推理的真正瓶颈。

5.3 精度对齐问题的处理

选型后模型结果偏了,经常是下面的原因:

TF32 模式问题。Tensor Core 在默认情况下可能用 TF32 替代 FP32,精度会掉一些。如果模型对精度敏感,需要显式关闭或改用 FP32 部分算子。

混合精度的梯度缩放问题。AMP 模式下 loss 突然不降或出现 NaN,多半是梯度溢出和 loss scaling 参数的问题。可以先跑 FP32 reference 对比,再逐层启用混合精度定位。

reduce 顺序不同。不同框架对矩阵乘内部的累加顺序不同,结果微小差异很正常。我判断是否可接受时,习惯计算余弦相似度而不是绝对误差。比如让输出的 logits 余弦相似度大于 0.99,基本可以认为行为对齐。

如果你用的是量化模型,INT8/INT4 的算子实现更敏感。AWQ、GPTQ 对量化顺序、缩放因子处理方式都会影响精度。选型前先拿你那几层敏感算子,单独对比量化前后的精度衰减。

5.4 选型避坑速查表

做一个速查表,方便你对照检查:

现象可能原因排查方法解决建议
模型无法编译算子目标框架不支持查看报错堆栈定位算子升级框架/替换等价算子/自定义
性能远低于预期未开融合、精度布局不对profiler 看 kernel 类型开启 torch.compile/CUDA Graph/混合精度
长请求吞吐上不去KV cache 带宽瓶颈监控 HBM 带宽利用率使用带分页缓存和量化缓存的推理框架
训练 loss 异常混合精度梯度溢出对比 FP32 baseline调整 loss scaling,逐层定位
多卡扩展性差通信算子耗时占比高profiler 看 AllReduce 时间换通信库,开启通信计算重叠
边缘端模型太大算子显存占用高检查算子显存分配用 INT8 量化、算子融合、分块推理

这些坑我基本都踩过。选型时提前对照这张表,能省很多不必要的加班。

最后再说一句我个人的习惯。每次选型我都会在测试环境单独建一个专门记录算子行为的脚本库,把每个候选框架在不同数据形状下的算子性能都沉淀下来。下次新模型来了,直接查历史数据,大部分情况下不用重新测。这个习惯帮我省了很多重复工作。选型不是一次性交付,而是一个可复用的过程资产,做好沉淀,后续每次模型升级都能快人一步。

RELATED READING

延伸阅读

更多一线实战笔记与深度复盘,助您持续精进