DeepGEMM：DeepSeek 为什么要自己写一套 BLAS 内核

项目：deepseek-ai/DeepGEMM · 363 stars/day · CUDA

一个反直觉的选择

NVIDIA 提供了 cuBLAS——经过十几年优化的 BLAS 库，几乎支持所有矩阵形状和数据类型。CUTLASS 更进一步，提供了可组合的 GEMM 模板。开源生态有 Triton、FlashAttention、xFormers。为什么 DeepSeek 还要自己写一套 GEMM 内核？

答案藏在一个具体的数字里：H800 上 1550 TFLOPS。这是 DeepGEMM 在 FP8 GEMM 上的峰值性能，接近 H800 理论 FP8 算力上限的 95%。

cuBLAS 做不到这个数字。不是因为 cuBLAS 不够好，而是因为 cuBLAS 是通用库——它要支持所有形状、所有数据类型、所有 GPU 架构。而 DeepSeek 只需要支持自己的模型用到的那些形状和数据类型。通用性的代价是性能，DeepGEMM 用放弃通用性换来了极致性能。

但这个故事比"自研更高效"复杂得多。

不只是 GEMM

DeepGEMM 的名字里有"GEMM"，但它已经不只是 GEMM 了。2026 年 9 月的更新让它变成了一个覆盖现代 LLM 核心计算原语的统一库：

注意最后一行——HyperConnection。这是 DeepSeek V4 架构特有的层间连接机制，不是标准 attention，不是标准 MLP，是 DeepSeek 自己设计的。没有现成的库支持它，因为它是全新的计算原语。

这揭示了一个结构性问题：当模型架构创新快于库的更新速度时，自研内核不是优化选择，是唯一选择。 cuBLAS 不会为你的新架构单独写一个 kernel——除非你是 NVIDIA 最重要的客户。CUTLASS 提供模板，但模板组合复杂度很高，且不一定能覆盖非标准操作。

DeepSeek 的选择是：自己写。而且写完之后开源。

DeepJIT：安装时不编译，运行时编译

DeepGEMM 最有意思的工程选择是 DeepJIT——运行时 JIT 编译。

传统 CUDA 库的流程：写 CUDA 代码 → 用 NVCC 编译成 .so 或 .cubin → 用户安装时拿到预编译二进制 → 运行时直接加载。问题是：不同 GPU 架构（SM80、SM90、SM100）需要不同的编译产物，不同 CUDA 版本之间有兼容性问题，用户安装时经常遇到"编译失败"或"架构不匹配"。

DeepGEMM 的做法：安装时不编译任何 CUDA 代码。 所有内核在运行时通过 DeepJIT 动态编译。用户 pip install deep_gemm，不需要装 CUDA toolkit，不需要处理 NVCC 版本冲突，不需要担心架构不匹配——JIT 会在运行时检测 GPU 架构，生成对应的 PTX，然后编译加载。

这不是新概念——Triton 做了类似的事。但 DeepGEMM 的 JIT 有一个不同：它生成的不是 Triton IR，而是直接的 CUDA 代码。这意味着它可以利用 CUTLASS 和 CuTe 的概念（TMA、wgmma、tcgen05），但不需要依赖它们的模板系统。

类比一下：Triton 像是 Python——你写高层代码，编译器帮你优化。DeepJIT 像是 C++ 模板——你写接近硬件的代码，编译器在运行时实例化。前者更易用，后者更灵活。DeepSeek 选择了后者，因为他们需要榨干最后一点性能。

简单性作为设计目标

DeepGEMM 的 README 里有一句话容易被忽略：

The library is designed for simplicity, with only a limited number of core kernel functions, making it a clean and accessible resource for learning NVIDIA GPU kernel optimization techniques.

"为简单性设计"——这听起来不像一个追求极致性能的库会说的话。但仔细想想，这是有道理的。

DeepGEMM 只有有限的几个核心 kernel 函数。每个函数对应一个明确的计算原语（dense GEMM、grouped GEMM、masked GEMM、MoE GEMM）。没有复杂的模板层级，没有需要理解十层抽象才能修改的代码结构。

这和 CUTLASS 形成对比。CUTLASS 是一个强大的库，但它的模板层级之深让很多人望而却步——修改一个 kernel 需要理解 TiledMma、ThrMma、SmemCopy、R2S、S2R 等一堆概念。DeepGEMM 把这些概念简化了，只保留必要的部分。

简单性不只是为了易用，也是为了可维护性。 DeepSeek 需要快速迭代——当新的模型架构（如 V4 的 HyperConnection）需要新的 kernel 时，他们需要能快速实现和部署。如果库本身太复杂，每次添加新原语都需要修改十层模板，迭代速度会慢到不可接受。

这里有一个权衡：通用性 vs 性能 vs 简单性。你最多只能选两个。cuBLAS 选了通用性+性能（牺牲简单性，因为闭源）。CUTLASS 选了通用性+简单性（牺牲一些性能，因为模板开销）。DeepGEMM 选了性能+简单性（牺牲通用性，只支持特定形状和架构）。

DeepSeek 能做这个选择，是因为他们同时控制了模型架构和推理框架。模型架构决定了需要哪些计算原语，推理框架决定了如何调用这些原语。当两者都在自己手里时，通用性就不再是必需品——你不需要支持所有形状，只需要支持你的模型用到的那些形状。

Ascend 支持：走出 NVIDIA

2026 年 9 月 30 日的更新：DeepGEMM-Ascend 发布。

这意味着 DeepGEMM 不再只是 NVIDIA 的库——它开始支持华为的 Ascend 芯片。这不是简单的移植，而是架构层面的扩展：Ascend 的达芬奇架构和 NVIDIA 的 Tensor Core 是完全不同的编程模型。

为什么这件事重要？因为它意味着 DeepSeek 在认真考虑 NVIDIA 替代方案。不是嘴上说说，而是在最底层的计算库层面做准备。

如果 DeepSeek V5 要在 Ascend 上推理，需要的不是把模型权重搬过去——需要的是整个计算栈的适配。从 GEMM 到 MoE 到注意力机制，每一个原语都需要在 Ascend 上有高性能实现。DeepGEMM-Ascend 是这个适配的第一步。

这也解释了为什么 DeepJIT 的设计是明智的。如果 DeepGEMM 是预编译的 NVIDIA 二进制，扩展到 Ascend 需要重写整个编译流程。但 JIT 架构只需要添加一个新的代码生成后端——运行时编译的灵活性在跨架构扩展时变成了真正的优势。

这意味着什么

DeepGEMM 表面上是一个 GEMM 库，实际上是一个垂直整合策略的底层拼图。

DeepSeek 控制了整个栈：模型架构（V3/V4 的设计）→ 训练框架 → 推理框架（DeepSeek 推理引擎）→ 计算内核（DeepGEMM）→ 硬件适配（从 NVIDIA 到 Ascend）。每一层都为其他层优化，没有"等上游库支持"的依赖。

这种垂直整合的代价是工程投入——维护一套自己的 GEMM 库需要持续的 GPU 内核专家投入。但收益是：没有人能卡你的脖子。NVIDIA 不更新 cuBLAS？无所谓。CUTLASS 不支持你的新架构？无所谓。Ascend 需要支持？自己加。

当大多数 AI 公司还在依赖第三方库做推理时，DeepSeek 已经把栈挖到了 GPU 内核层面。这不是过度工程——当你每月的推理账单以百万美元计时，5% 的性能差异就是真金白银。而 DeepSeek 的推理成本，正是他们能以十分之一的价格提供 API 的核心原因之一。

DeepGEMM 开源的策略也值得玩味。开源一个自己深度依赖的库，意味着社区可以帮你发现 bug、贡献优化、扩展到更多硬件。而核心的模型架构和训练数据仍然闭源。开源基础设施，闭源上层价值——这个策略 DeepSeek 不是第一个用的，但他们在内核层面的执行深度是罕见的。

下一步值得关注的不是 DeepGEMM 能不能再快 5%，而是当 Ascend 支持成熟后，DeepSeek 会不会在 Ascend 集群上跑推理。如果那一天到来，NVIDIA 在 AI 推理市场的定价权就会真正受到挑战——不是来自政策制裁，而是来自技术替代。