大模型竞赛中，算力不再只是堆显卡，更是抢效率。

面对 H20 等推理卡在主流算子库下难以跑满性能的痛点，腾讯混元 AI Infra 团队正式开源生产级高性能 LLM 推理核心算子库 HPC-Ops。

该算子库采用 CUDA 和 CuTe 从零构建，通过抽象化工程架构、微架构深度适配及指令级极致优化等，降低底层算子开发门槛，将核心算子性能逼近硬件峰值，实现了显著性能突破。

在真实场景下，基于 HPC-Ops，混元模型推理 QPM 提升30%，DeepSeek 模型 QPM 提升17%。

{jz:field.toptypename/}

同时，在单算子性能方面，HPC-Ops 实现 Attention 相比 FlashInfer/FlashAttention 最高提升 2.22 倍；

GroupGEMM 相比 DeepGEMM 最高提升 1.88 倍；FusedMoE 相比 TensorRT-LLM 最高提升 1.49 倍。

主流算子库亟需更适配的底层支持

在大模型时代，计算效率已成为 AI 应用及发展的关键瓶颈。

目前主流算子库（如 FlashInfer、DeepGEMM）多以NVIDIA H800等高配训练卡为首要优化目标，但限于客观原因，不少大模型的大规模线上推理服务只能采用 H20 等推理型计算卡。

现有 SOTA 算子库在这些显卡上往往难以发挥硬件峰值能力。

同时业务侧对极致吞吐、低延迟以及 Blockwise FP8 等复杂量化策略的需求日益迫切，亟需更适配的底层支持。

总结来看，现有主流算子库主要存在以下痛点。

使用成本高

主流算子库设计复杂，核心 Kernel 封装深，在其上修改适配成本非常高，除了对代码非常熟悉的开发者，普通的 AI 研究者很难在其上适配修改。

而大模型的很多加速研究创新，比如量化算法和投机采样等方法都严重依赖于与之匹配的高效算子实现。

比如最开始 4bit 和 8bit 的量化算法出来后，虽然理论上加载数据量减少，但由于没有与之匹配的低精度算子实现，低精度量化在很长的一段时间内都是负优化。

目标硬件不匹配

现有的主流算子库都是以 H800 等显卡为目标优化、NVIDIA 提供的 CUTLASS 等算子更是以 Blackwell 架构为目标，而目前国内主流的推理显卡则有所不同。

不同硬件间算力带宽的差距导致 Kernel 的优化方法也会不同，因此现有的算子库在国内主流推理卡上的表现并未发挥出硬件的全部性能。

基于以上问题，腾讯混元使用 CUDA 和 CuTe 开发了一套轻量、高效的 LLM 核心算子库。

用 CUDA 和 CuTe 从零构建

该算子库主要包括 FusedMoE、Attention、机内 / 机间通信、Norm、Sampler、以及各类小算子的融合算子模块，整体算子库架构如下图所示。

通过分析任务特性和硬件微架构，将任务的划分逻辑与硬件指令做了更好的对齐，以此获得更好的性能，并且对工程代码进行了适度的抽象，让开发者能聚焦于算法本身，降低维护门槛。

该算子库不仅是高性能生产工具，也可作为开发者深入理解 CUTLASS 与 CuTe 工业级开发的实践范本，具体的技术细节如下。

任务特性与硬件能力对齐

针对访存瓶颈的算子，其性能主要受限于数据加载速度。

针对国内的主流推理显卡，通过调整指令发射顺序进行数据预取优化，确保数据传输单元一直处于高利用率。

针对不同的问题规格做了更细致的指令对齐和优化，去除冗余低效指令以减少算力的浪费，如针对 Decode Attention 和小 batch 下的 GroupGEMM 都做了 AB 矩阵交换的优化；

以此对齐到硬件架构上的 wgmma 指令，访存带宽可达到硬件峰值能力的 80% 以上。

Attention SwapAB 示意图精细的任务调度和数据重排

针对每个算子问题，都重新思考了任务数据的划分调度策略，尽可能保证每个 SM 都任务均衡的同时兼顾 cache 的连续性。

并且采用了 persistent kernel 的方式隐藏 kernel prologue 和 epilogue 的开销。

另外也通过数据重排减少了额外的操作和显存占用。

比如在 FP8 Attention Kernel 中创新性采用了 Interleave 重排技术，解决了指令不匹配的问题，开云中国app登录入口减少线程间数据 shuffle，获得了优于业界 SOTA 的算子性能。

聚焦于计算逻辑本身

GPU 编程的复杂度很大程度上来源于操作的复杂性，为了能使用高效指令，一般需要对数据进行多次的重解释和变换等编程技巧，这大大加重了开发者的心智负担。

因此基于 CuTe 扩展开发 vec 抽象层统一负责高效数据搬运，利用 Layout 代数抽象隔离复杂的 Tiling 与计算逻辑，让开发者能聚焦于算法本身，降低维护门槛。

关键实验结果

通过以上高效算子实现，在混元模型上将 QPM 端到端提升 30%，DeepSeek 上 QPM 提升 17%。

同时针对 LLM 中核算子模块进行了测试，以常用的模型规格（混元、DeepSeek）进行了测试，并对比了目前主流的算子库实现。

实验表明，在 LLM 的核心模块 Attention 和 FusedMoE 上的性能都超越当下 SOTA 实现。

GroupGEMM

与 DeepGEMM ( v2.2.0 ) 的两种版本进行对比，在 Batch

且通过流水线掩盖技术 Blockwise 与 PerTensor 性能几乎持平；

在大 Batch 场景下，亦能保持约 1.1x 的领先优势。该算子同时兼容紧密排布与 Token 不连续输入，显著减少临时显存用量。

GroupGEMM 性能对比图 FusedMoE

完整封装了包括前序数据重排、GroupGEMM 及后续 Reduce 加权平均在内的全流程模块 .

并在序列长度取 16 倍数的均衡分配规格下，对比了 vLLM （v0.11.0）与 TensorRT-LLM （v1.1.0）的实现。

测试结果显示，该 FusedMoE 模块在 TP 场景下相比 TensorRT-LLM 提升显著，最大性能提升达 1.49x；在 EP 模拟均衡场景下最大提升 1.09x。

针对不同输入长度采取的差异化重排策略，进一步确保了整体模块在各类规格下的最优表现。

FusedMoE 性能对比图 Attention

针对 Prefill 场景，测试 128～64K 的输入长度。

在 batch 较小时，BF16 精度下相比 SOTA 实现提升 1.3x；在大 batch 时基本与当前 SOTA 对齐。

针对 Decode 场景，根据线上 SLO 约束，搭配一组 batch 和输入长度的测试用例，BF16 精度下提 1.35x～2.22x；

FP8 精度下，当 Sequence Length 较小时与 SOTA 相当，当 Sequence Length 较大时相比 SOTA 提升 1.09x~2.0x。

Attention 性能对比图算子库当前能力和未来发展方向

作为面向大模型推理场景的高性能算子库，HPC-Ops 凭借 Attention、FusedMoE、GroupGEMM 等核心算子的极致优化，达成最高 2.22 倍的性能提升，且已在腾讯大规模生产环境中完成验证。

其简洁易用的 API 可无缝对接 vLLM、SGLang 等主流推理框架，原生支持 BF16、FP8 等多精度量化方案。

同时还以 CuTe、CUTLASS 为基础，提供了数百行代码即可构建 SOTA 算子的实践范例，为开发者降低了高性能 CUDA 内核的开发门槛。

在未来的发展规划中，HPC-Ops 将持续深耕大模型推理性能的突破方向。

一方面，将重点研发稀疏 Attention 算子，针对性解决长上下文大模型的内存与算力瓶颈；

另一方面，会拓展更丰富的量化策略，覆盖 4bit/8bit 混合精度等更多量化方案，进一步平衡推理速度与模型精度。

此外，算子库还将布局计算 - 通信协同优化的内核，通过融合多 GPU 间的计算逻辑与通信流程，大幅降低分布式推理场景下的通信开销，为超大规模大模型的高效部署提供底层支撑。

目前，HPC-Ops 已在 GitHub 开放源码供开发者下载使用。

同时腾讯混元 Infra 团队也表示，欢迎行业内的技术实践者提交高价值 PR，参与算子边缘场景优化、教程案例打磨等精准化贡献，共同推动大模型推理技术的边界拓展。

GitHub 项目地址：https://github.com/Tencent/hpc-ops

一键三连「点赞」「转发」「小心心」

欢迎在评论区留下你的想法！

—  完  —

我们正在招聘一名眼疾手快、关注 AI 的学术编辑实习生  � �

感兴趣的小伙伴欢迎关注 � �  了解详情

� � 点亮星标 � �

科技前沿进展每日见