【免费下载链接】hpc-opsHigh Performance LLM Inference Operator Library项目地址https://gitcode.com/gh_mirrors/hp/hpc-ops点击查看免费下载hpc-ops 是腾讯混元 AI Infra 团队开源的 LLM 推理高性能 CUDA 算子库High Performance LLM Inference Operator Library聚焦 Attention、Fused MoE、GEMM、采样器、归一化与通信-计算融合这几条决定真实推理延迟的热路径。在 NVIDIA H20SM90上其 decode Attention 最高领先 FlashInfer/FlashAttention-3 等基线2.22 倍融合 MoE 相比 vLLM、SGLang 最高领先1.6 倍且已在腾讯大规模生产推理中验证。为什么 LLM 推理需要一个新的 CUDA 算子库 大模型在线服务的延迟瓶颈往往不在模型本身而在底层算子的工程细节Decode 阶段每个 token 只算一次 Attention但 KV Cache 长度差异巨大静态调度会产生长尾延迟MoE 路由gather → GEMM → 激活 → GEMM → 归约的链条中每一步的 kernel launch 和 HBM 读写都被放大张量并行AllReduce 残差加法 RMSNorm 本是同一个逻辑操作拆成三段却要三次 launch、多次读写显存。hpc-ops 的做法是把每条链路融进尽可能少的 kernel并用 TMA、cp.async、PDLProgrammatic Dependent Launch、CUDA multicast 等 SM90 新特性做深度优化。对普通用户来说最大的好处是Python API 干净可无缝替换进 vLLM / SGLang 等框架见 hpc/ 下的封装文件。算子目录一览覆盖推理热路径的每一环领域算子主要场景精度Python 入口AttentionPrefill / Decodepaged KV Cache动态 decode 调度变长请求的在线服务BF16 / FP8hpc/attention.pyAttentionFP8 块稀疏 Prefill长上下文 块级稀疏 maskFP8hpc/stem.pyGEMMBF16 × FP32 GEMMMoE 路由 GEMM、状态压缩 GEMMBF16×FP32hpc/gemm.pyGrouped GEMMFP8 GroupGEMM专家并行的分组矩阵乘FP8hpc/group_gemm.pyFused MoE融合 FP8 MoE 流水线TP/EP 下的低延迟 MoEFP8hpc/fuse_moe.py通信AllReduce 残差 RMSNorm 融合张量并行后处理BF16hpc/allreduce.py采样融合 Samplerdecode 步的 token 采样BF16 / FP32hpc/sampler.py工具类RMSNorm / RoPEKV 写入 / 激活量化通用 LLM 层工具BF16 / FP8hpc/rope.py、hpc/normalization.py核心算子深度解析 1. Attention从静态 split-k 到动态任务调度在线 decode 的 batch 里各请求的 KV 长度可能相差一个数量级例如1×128K 31×4K这样的极端混合负载。静态 split-k 无法适配这种分布长请求拖出 CTA 级长尾短请求却浪费算力。hpc-ops 的动态 decode Attention入口见 hpc/attention.py 中的assign_attention_decode_task把每个请求统一切分为等长 KV tile在每一步 decode 前按贪心装箱策略把 tile 均衡分配给各 CTAkernel 直接消费任务图再由 combine kernel 合并 split-k 结果。混合长度场景下最高带来2.88 倍加速。长上下文场景还有FP8 块稀疏 Prefill Attention接收预计算的块 mask整块跳过无关 KV tile同时用 per-tile FP8 缩放保证数值质量最高3.16 倍加速。源码位于 src/attention/prefill/sm90/warp_spec_with_kvcache_blocksparse_fp8_dim128.cudecode 侧在 src/attention/decode/sm90/dynamic/。2. Fused MoE一条流水线吃掉路由到归约的全部环节FusedMoE 路径hpc/fuse_moe.py把路由 → Gate-Up GEMM → 激活量化 → Down GEMM → top-k 加权和归约融进一条流水线Gate-Up GEMM 通过路由索引直接读原始 token省掉独立的 gather 阶段及其额外显存流量用 shared memory 计数 连续 expert 输出区间降低全局 atomic 压力低延迟 regime 下用 cp.async 路径 去掉 Warp Specialization提高 CTA 驻留率把延迟隐藏从CTA 内软件流水线转移到跨 CTA 硬件调度再用 PDL 把各阶段串起来消除 kernel launch 空隙。在 DeepSeek-V3、Hunyuan-V3、Qwen3-235B 三种模型形状下TP8/EP1 与 TP1/EP8 两种部署形态中最高领先 vLLM CUTLASS、vLLM Triton、SGLang1.6 倍TP/ 1.5 倍EP复现脚本见 benchmark/fused_moe/README.md。3. GEMM 家族BF16×FP32 的精度-速度两全解现代 Tensor Core 没有原生 FP32 吞吐朴素 FP32 GEMM 会退回 CUDA Core而直接用 BF16/TF32 又会损失精度。MoE 路由 GEMM 恰好属于激活是 BF16、权重需 FP32 精度的场景。hpc-ops 的 BF16×FP32 GEMM 思路很巧把每个 FP32 权重拆成高 BF16 分量 低 BF16 残差固定缩放 1/256再在单个 kernel内完成两个 BF16 Tensor Core GEMM 的融合线性组合——共享输入搬运、累加器留在寄存器、最终结果只写一次。相比 cuBLAS FP32/TF32 最高3.22 倍精度对齐 FP32 级。GroupGEMM FP8hpc/group_gemm.py则面向专家并行的分组矩阵乘decode 阶段最高领先 DeepGEMM1.88 倍。4. 通信-计算融合AllReduce 残差 RMSNorm张量并行里RMSNorm(AllReduce(x) residual)被切成三段执行时launch 开销和 HBM 读写都会被放大。hpc-ops 把它融成NVLink 原生 kernelhpc/allreduce.py高吞吐模式fuse_allreduce_rmsnorm_high_throughput大 token 的 prefill 形状用 CUDA multicast 做 two-shot 通信低延迟模式fuse_allreduce_rmsnorm_low_latency小 token 的 decode 形状Lamport P2P 双 kernel 设计 PDL 重叠。对比 NCCL、FlashInfer 最高1.76 倍。5. 融合 Sampler小 batch 下的隐藏杀手decode 采样链路重复惩罚 → 温度缩放 → top-k → top-p → softmax → 随机采样 → 惩罚 mask 更新通常是一串小 kernel 对 vocab 维度的反复扫描。hpc-ops 的 融合采样器 把整条链路压到两个 CUDA kernel温度-only 场景还有更轻的 fast pathvocab 大小 120832、batch 1~512 时最高领先 vLLM 风格 PyTorch / FlashInfer8.5 倍。性能速览数据说话 算子对比基线代表性加速Attention BF16prefill/decodeFlashInfer、FA2、FA3、TensorRT-LLM最高 1.33x /2.22xAttention FP8prefill/decodeFlashInfer、FA3、TensorRT-LLM最高 1.12x / 2.0x动态 decode 调度静态 split-k最高 2.88xFP8 块稀疏 PrefillMIT-BSA、FlashPrefill-BSA、FA3最高 3.16xBF16×FP32 GEMMcuBLAS FP32 / TF32最高 3.22xFused FP8 MoEvLLM CUTLASS/Triton、SGLang最高 1.6x / 1.5xAllReduceRMSNorm 融合NCCL、FlashInfer最高 1.76x融合 SamplervLLM 风格 PyTorch、FlashInfer最高 8.5xGroupGEMM FP8DeepGEMM最高 1.1x / 1.88x性能因场景而异所有数字均可用仓库自带 benchmark 复现。一键安装步骤 ⚡环境要求NVIDIA SM90 架构 GPUH20/H100、Python ≥ 3.8、支持 C17 的编译器、CUDA ≥ 12.9。git clone https://gitcode.com/gh_mirrors/hp/hpc-ops cd hpc-ops # 构建覆盖所有支持架构的 wheel make wheel python3 -m pip install dist/*.whlwheel 内部按 SM 架构各带一个扩展模块import hpc时 hpc/__init__.py 会自动探测本机 GPU 并加载匹配模块——sm90 的模块不会跑在 sm103 设备上加载即报错避免在 kernel 深处崩溃。最小使用示例FP8 GroupGEMMimport torch import hpc x torch.randn(1024, 4096, devicecuda).to(torch.float8_e4m3fn) w torch.randn(8, 4096, 4096, devicecuda).to(torch.float8_e4m3fn) output hpc.group_gemm_pertensor_fp8(x, w, num_tokens_per_group, cu_num_tokens_per_group, scale)其余算子用法可直接参考 tests/ 下对应的测试文件如 tests/test_fuse_moe_pertensor.py、tests/test_attention_decode_bf16.py。如何复现 Benchmark新手友好指南每个算子都有独立的基准目录全部支持nsys CUDA Graph replay 的发布级计时方式Fused MoE对比 vLLM / SGLangbenchmark/fused_moe/benchmark_fuse_moe.pyDecode Attention对比 FlashInfer / FA3benchmark/attention_decode/Route GEMM对比 cuBLASbenchmark/route_gemm/benchmark_gemm_bf16xfp32.py融合 AllReduceRMSNorm对比 NCCLbenchmark/fuse_allreduce_rmsorm/Sampler、RoPEKV 写入benchmark/sampler/、benchmark/rope_norm_store_kv/例如复现 MoE 对比python3 benchmark/fused_moe/benchmark_fuse_moe.py \ --tp 8 --ep 1 \ --providers hpcops vllm vllm_cutlass sglang \ --csv fused_moe_tp8_ep1.csvvLLM / SGLang 属于可选对比项本地没装会自动跳过并继续不会中断整个 sweep新手环境也能顺利跑通。代码怎么读给想学 SOTA CUDA 的新手 hpc-ops 官方定位为一份实战级现代 CUDA 教程目录规则极简看路径就知道一切Python 层hpc/ 下一个文件一个算子族定义用户 API 与 fake 实现Host 层每个算子的 entry.cc 负责校验、分配输出、选 launch 配置并用HPC_ARCH_DISPATCH宏在 launch 点分发到具体架构Kernel 层smarch/目录下的.cu/.cuh就是该架构的实现比如 src/fuse_moe/sm90/fuse_moe.cuCUTLASS/CuTe第三方 CUTLASS 以 vendored 方式放在 3rd/cutlass/TMA、cluster 等基础设施可随源码一起读。想深入细节开发规范都写在 CLAUDE.md 里——从源码选择规则到跨算子调用约定是一份少见的高质量仓库地图。总结与 Roadmap hpc-ops 证明了贴着真实 serving 负载打磨的 CUDA kernel可以在 Attention、MoE、GEMM、采样、通信融合这些环节全面超越通用框架自带实现而且 API 足够干净普通用户几行代码即可接入。官方后续路线包括 4/8bit 混合精度扩展、Megakernel 多算子单 kernel 化、更多 GPU 架构支持以及低精度 AllReduce/AllGather 通信 kernel。如果你是 LLM 推理工程师建议从 benchmark/ 复现数字开始如果想学现代 CUDA直接读 src/ 下 TMA、PDL、multicast 的生产级用法会比任何教程都更快。赞分享【免费下载链接】hpc-opsHigh Performance LLM Inference Operator Library项目地址https://gitcode.com/gh_mirrors/hp/hpc-ops点击查看免费下载相关推荐如何在Windows上免费获得苹果平方字体的终极指南 如何在Windows上免费获得苹果平方字体的终极指南 你是否曾经羡慕Mac电脑上那些清晰锐利的中文字体是否在为Windows系统单调的字体选择而苦恼别前端django-friendship用户指南掌握好友请求、关注和屏蔽功能的完整攻略django friendship用户指南掌握好友请求、关注和屏蔽功能的完整攻略 django friendship是一个功能强大的Django应用专门用于Boss Show Time5大核心功能帮你精准把握招聘黄金时间Boss Show Time5大核心功能帮你精准把握招聘黄金时间 在当今竞争激烈的求职市场中时间就是机会。Boss Show Time作为一款智能时间展示插前端创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考