2026-09-30，DeepSeek 正式开源面向华为昇腾算力平台的基础设施组件，一口气放出了四个可以直接进代码仓库的项目：算子库 TileKernels、矩阵计算库 DeepGEMM-Ascend、分布式通信库 DeepEP-Ascend，以及注意力算子库 FlashMLA 的昇腾版本。按官方口径（经 AI工具集每日快讯 09-30 期转述），这批组件与此前英伟达平台上的 DeepGEMM、DeepEP、FlashMLA 一一对应，覆盖矩阵运算、通信、稀疏注意力等核心能力，多项测试接近硬件上限。
对大多数工程师来说，真正的门槛不是「想不想学」，而是手上有没有一块 Ascend 950。这篇 SOP 就按门槛把上手路线拆成三条：没有昇腾硬件，走读码路线，照样能吃透算法与工程取舍；有昇腾卡，从 pip 安装 TileKernels 开始跑算子；要搭多卡专家并行，再上 DeepEP-Ascend，同时把它当前的许可证状态与验证栈边界搞清楚。全文的命令与版本号都取自各仓库 README 原文，不掺编造成分。往下看之前，先做个分工说明：本站 DeepSeek V4.1 Flash 开源热点讲模型本身的开源情况，DeepSeek V4.1 Flash 生产接入 SOP讲 API 层怎么接入生产，本文只管算子栈这一层，三层互为补充，不互相重复。
算子栈里有什么：四个仓库各管一段
先建立全局认知。这四个仓库的定位并不重叠，可以按「算子、矩阵、通信、注意力」四层来记：
- TileKernels：基于 TileLang（支持多硬件后端的领域专用语言）实现的数十个深度优化算子库，覆盖 MoE 路由（top-k 专家选择）、量化（per-token/per-block/per-channel 的 FP8/FP4 cast 与反量化、融合 SwiGLU 加量化）、Engram 门控（融合 RMSNorm、前反向与权重梯度归约）、流形超连接 mHC（含 Sinkhorn 归一化）、RoPE 与随机数等。README 官方自报，大多数算子性能接近硬件的计算吞吐或带宽上限，且全部算子已用于 DeepSeek 内部训练与推理。2026-09-30 起加入华为昇腾支持：沿 NVIDIA 路线新增第二后端，运行时自动选择，同一套 Python API 可同时跑 NVIDIA GPU 与昇腾 NPU。
- DeepGEMM-Ascend：DeepGEMM 的昇腾移植，与 DeepGEMM 完全 API 兼容，连包名都是同一个 deep_gemm，支持 BF16、FP8、FP4 GEMM、MQA logits 与 MegaMoE 算子。它对昇腾 MAD（矩阵乘加）原语做了轻量抽象，隐藏分形布局、对齐约束与地址计算，并采用稀疏数据加载、协程流水线等昇腾特有优化。2026.09.30 首发即支持 Ascend 950 设备。
- DeepEP-Ascend：面向昇腾 NPU 的高性能训练与推理通信库，提供 MoE dispatch/combine 的专家并行 all-to-all 操作，支持 FP8 dispatch 与延迟 epilogue，另有流水线并行、Bucket 集合通信与 Engram 远端内存访问等通信原语（部分在开发中）。其公开 buffer API 与 NVIDIA 版 DeepEP 对齐，训练、推理 prefill 与 decoding 共用同一 EPBuffer 接口和 deep_ep 包名，实现遵循 DeepEP 的 EPBuffer V2.5 API。
- FlashMLA：DeepSeek 的高性能注意力算子库，支撑 DeepSeek-V4.1 模型在 NVIDIA GPU 与华为昇腾 NPU 上的推理。2026.09.30 发布昇腾侧稀疏注意力 prefill 与 decoding 算子，按 README 官方口径，prefill 达 410 TFlops（95% 硬件峰值），decoding 达 360 TFlops（83% 硬件峰值），并附算法与优化技术深度报告。
英伟达侧的对应主仓 DeepGEMM、DeepEP、FlashMLA 均为 MIT 许可证，这为对照阅读提供了完整参照系。国产芯片训练这条线，此前本站 美团 LongCat 2 热点里只有美团公开走国产芯片训练路线，DeepSeek 这次直接把整套算子栈摆上了台面，含金量不是一个量级。
门槛总表：哪些要上机，哪些读码就够
三条路线的硬件与软件门槛差异很大，先看总表再决定走哪条：
| 路线 | 硬件门槛 | 软件门槛 | 你能得到什么 |
|---|---|---|---|
| 读码路线 | 无需昇腾硬件，能上网即可 | 无 | 算子实现细节、移植差异、深度技术报告 |
| TileKernels 与 DeepGEMM-Ascend | Ascend 950 NPU（TileKernels）；DeepGEMM-Ascend 在 Ascend 950 系列开发验证 | TileKernels 要求 Python 3.12+、PyTorch 2.13+、TileLang 0.1.15+、CANN 9.2.0+；DeepGEMM-Ascend 要求 CANN 9.20、torch_npu、Python 3.10+、支持 C++20 format 的编译器 | 跑通量化、路由、GEMM 等单卡算子 |
| DeepEP-Ascend | Ascend 950（A5），要求 UBMEM 连通性与 UBC_CTP/URMA 通道，多卡通信需偶数个 rank | 验证栈为 Ascend 950DT、CANN 9.2.0、Python 3.12、PyTorch 2.13.0+cpu、torch_npu 2.13.0rc1，另需 PoC HDK 配置 | 专家并行 all-to-all 通信 |
表里有个值得注意的分层：TileKernels 是唯一提供 pip 包的仓库，昇腾后端要求明确写着 Ascend 950 NPU 加 CANN 9.2.0 及以上；DeepGEMM-Ascend 目前只能源码安装；DeepEP-Ascend 的门槛最高，因为它管的是卡与卡之间的通信。CUDA 用户则完全不必等昇腾硬件：TileKernels 的 CUDA 后端要求 SM90 或 SM100 架构 GPU 加 CUDA 13.1 及以上，英伟达侧主仓直接可用。
路线一：没有昇腾硬件的读码路线
没有昇腾卡丝毫不妨碍学这套东西，因为算子实现、性能数据、深度报告全部公开。建议按这个顺序推进。
第一步，把四个仓库克隆下来。注意其中两个带子模块，克隆方式有讲究（后面「常见坑」细说）：
# TileKernels：TileLang 算子库
git clone https://github.com/deepseek-ai/TileKernels.git
# DeepGEMM-Ascend：子模块必须一并克隆
git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git
# DeepEP-Ascend：通信库
git clone https://github.com/deepseek-ai/DeepEP-Ascend.git
# FlashMLA：注意力算子，同样带子模块
git clone https://github.com/deepseek-ai/FlashMLA.git flash-mla
cd flash-mla
git submodule update --init --recursive第二步，读 FlashMLA 的昇腾深度技术报告。仓库内文档路径为 docs/20260930-ascend-prefill-deep-dive.md（中英双版），内容是昇腾 prefill 算子背后的算法与优化技术。这是本次开源里少见的「直接给方法论」的材料：它解释了稀疏注意力 prefill 在昇腾 950 上怎么打到 95% 硬件峰值，比只看代码省力气得多。
第三步，对照英伟达侧主仓看移植差异。英伟达侧的 DeepGEMM、DeepEP、FlashMLA 都是成熟仓，两相对照，移植思路一目了然。几个值得专门留意的差异点：DeepGEMM-Ascend 对昇腾 MAD 原语做了轻量抽象，并用稀疏数据加载、协程流水线等昇腾特有优化；昇腾侧的 scaling factor 格式与英伟达不同（下文「常见坑」详述）；DeepEP-Ascend 的 Ascend C 内核走 HCCL/HCOMM、UBMEM 与 URMA 通信栈，并通过 DeepJIT 在运行时编译，这与 CUDA 侧的编译模型完全不同。
CUDA 用户可以直接用英伟达侧主仓，不用等昇腾侧。但要提醒一句：FlashMLA 里那个把 Q-norm、Q-RoPE、注意力、O-RoPE 与 FP8 cast 融合起来的 kernel，目前仅支持 CUDA，昇腾不支持，这是 README 原文明确写的。另外 FlashMLA 的 2026.09.30 版本有破坏性变更：移除了对 Hopper 架构与前代模型（DeepSeek V3、V3.2、V4.0）的支持，并改了 FP8/FP4 KV cache 格式，与旧版本不兼容；需要跑旧模型的要切到官方指明的旧 commit。
路线二：有昇腾硬件，先装 TileKernels 与 DeepGEMM-Ascend
有了昇腾 950，最顺的起点是 TileKernels，因为它直接发 pip 包。
# 安装发布版
pip install tile-kernels
# 或源码安装开发版
pip install -e ".[dev]"装完可以直接跑仓库自带的 pytest 验证。README 给出的测试方式如下，加 --run-benchmark 参数可同时跑正确性与基准测试：
python -m pytest tests/quant/test_per_token_cast.py -n 4
python -m pytest tests/quant/test_per_token_cast.py --run-benchmark环境要求再复述一遍：Python 3.12 及以上、PyTorch 2.13 及以上、TileLang 0.1.15 及以上；昇腾后端要求 Ascend 950 NPU 加 CANN 9.2.0 及以上。因为后端在运行时自动选择，同一套 Python API 换台 NVIDIA 机器（SM90 或 SM100、CUDA 13.1+）也能直接跑，这对写跨平台测试很有用。
接下来是 DeepGEMM-Ascend，目前需源码安装。先确认环境：CANN 9.20 工具链（要提供 bin/bisheng 与 bin/ld.lld）、torch_npu、Python 3.10 及以上、支持 C++20 format 的编译器；它还依赖 tilelang（HC prenorm kernel 会用到）。安装流程：
# 子模块必须一并克隆
git clone --recursive https://github.com/deepseek-ai/DeepGEMM-Ascend.git
cd DeepGEMM-Ascend
# 链接部分关键头文件并构建 C++ 扩展
cat develop.sh
./develop.sh
pip install . --no-build-isolation装好后用法与 DeepGEMM 完全一致：包名同为 deep_gemm，kernel 接口直接参考 DeepGEMM 的接口文档即可。它还提供 set_num_sms、set_npu_arch、transform_sf_into_required_layout 等工具函数，JIT 相关行为可通过 DG_JIT_DEBUG、DG_JIT_CACHE_DIR 等环境变量控制，缓存目录默认在 $HOME/.dj。
路线三：DeepEP-Ascend 通信库，连同合规注意一起看
DeepEP-Ascend 是门槛最高的一层，装之前先把三件事弄清楚。
第一，硬件与验证栈边界。它要求 Ascend 950（A5）具备 UBMEM 连通性与 UBC_CTP/URMA 通道，多卡通信要求偶数个 rank。README 明确写了验证栈：Ascend 950DT、CANN 9.2.0、Python 3.12、PyTorch 2.13.0+cpu、torch_npu 2.13.0rc1，外加手动配置的 PoC HDK。其他昇腾代际或 CANN 版本上的 kernel 支持未经这些测试确立，也就是说偏离这套栈就是自己开路。固件方面，华为 Q3 商用 HDK（Atlas 850E）公开时间按厂商发布计划预计在 2026 年 10 月中旬（约 10 月 15 日），届时应作为公开部署基线；README 现有性能数据是用交付给 DeepSeek 的 PoC HDK 手动配置测的，不是那个未发布的商用版的结果。
第二，许可证状态。必须如实说：截至 2026-10-01 快照，DeepEP-Ascend 仓库根目录没有 LICENSE 文件（GitHub API 的 license 字段为空，根目录文件列表中也无 LICENSE）。这与 DeepGEMM-Ascend、TileKernels、FlashMLA 三个 MIT 仓形成鲜明对比。对个人读码学习影响不大，但商用或对合规敏感的团队，建议保持观察、等官方补上许可证再引入，不要默认它是 MIT。
第三，安装流程。先激活 CANN 环境、装好配套的 PyTorch 与 torch_npu，设置 ASCEND_HOME_PATH 指向当前 CANN 安装（也可用 ASCEND_TOOLKIT_HOME），然后：
git clone https://github.com/deepseek-ai/DeepEP-Ascend.git
cd DeepEP-Ascend
# 通过 HTTPS 拉取所需的 DeepJIT 子模块
git submodule update --init --recursive third-party/deep_jit
python -m pip install --no-build-isolation .宿主侧扩展随安装构建，设备侧 kernel 首次使用时才编译，所以运行时要保持 CANN 工具链与 Bisheng 编译器可用。开发模式可用 bash develop.sh 把扩展构建并链接进源码树。想复现 README 的 EP 性能测试，先设好环境变量再跑测试脚本：
export EP_AVOID_RECORD_STREAM=1
export TASK_QUEUE_ENABLE=0
export HCCL_IF_BASE_PORT=48361
python tests/ep/test_ep.py --num-processes 8 --num-ai-cores 32 \
--num-tokens 16384 --hidden 7168 --num-topk 6 --num-experts 256 \
--dispatch-dtype fp8 --test-first-onlyREADME 的性能测试条件是每 rank 16384 tokens、hidden 7168、在 256 个专家里做 top-6 路由、32 AI cores 加 64 AIVs、FP8 dispatch 行主序 scale、BF16 combine、10 次 warmup 加 50 次采样，并注明持续 dispatch 带宽在 EP32 及以下约达物理载荷带宽上限的 90% 到 95%，更大 EP 与 combine 仍在优化中。具体数字建议直接看仓库图表，本文不转引。
常见坑：四个高频翻车点
第一坑，scaling factor 格式。昇腾侧与英伟达侧不同：沿 K 维度每两个 UE8M0 scaling factor 打包进一个 int16，且打包值按 MN-major 顺序存储，以获得最优硬件效率。从英伟达侧代码移植或迁移数据布局时，这里几乎必踩，DeepGEMM-Ascend 提供的 transform_sf_into_required_layout 工具函数就是为转换布局准备的。
第二坑，偶数 rank。DeepEP-Ascend 的多卡通信要求偶数个 rank，单测起 8 进程是惯例。如果你用奇数个 rank 起通信，属于 README 明确不支持的用法，别拿它当 bug 报。
第三坑，子模块没拉全。DeepGEMM-Ascend 要求 git clone --recursive；DeepEP-Ascend 依赖 DeepJIT 子模块，要单独执行 git submodule update --init --recursive third-party/deep_jit；FlashMLA 也要 submodule update --init --recursive。子模块缺失的典型症状是构建阶段报头文件或依赖找不到。
第四坑，--no-build-isolation 不能省。DeepEP-Ascend 安装命令原文带 python -m pip install --no-build-isolation .；DeepGEMM-Ascend 是 pip install . --no-build-isolation。这类 C++ 扩展项目在构建阶段要引用环境里的 torch 或 torch_npu，隔离构建环境会直接失败。
自查清单
动手前后照着这张清单过一遍：
- 环境核对：昇腾硬件型号（950 还是 950DT）、CANN 版本（9.2.0 还是 9.20）、Python、PyTorch、torch_npu 版本与目标仓库 README 要求一致。
- 子模块：DeepGEMM-Ascend 用 --recursive 克隆；DeepEP-Ascend 单独拉取 third-party/deep_jit；FlashMLA 执行 submodule update --init --recursive。
- 环境变量：ASCEND_HOME_PATH（或 ASCEND_TOOLKIT_HOME）已指向当前 CANN 安装；跑 DeepEP 测试前 EP_AVOID_RECORD_STREAM 已设置。
- 构建参数：安装命令带 --no-build-isolation；DeepGEMM-Ascend 先跑 develop.sh。
- 合规确认：DeepEP-Ascend 无 LICENSE 文件，商用前单独评估；其余三仓为 MIT。
- 版本预期：FlashMLA 2026.09.30 版与旧版不兼容（KV cache 格式变更），跑旧模型切旧 commit；DeepEP-Ascend 仅在验证栈内可信。
- 数据布局：昇腾侧 scaling factor 为 UE8M0 打包 int16、MN-major 存储，与英伟达不同，迁移前先转换。
常见问题
Q1： 没有昇腾硬件，这套算子栈能学到什么程度？
能学到七成以上。四个仓库的算子实现全部公开，FlashMLA 还附了昇腾 prefill 深度技术报告，把算法与优化思路讲透。英伟达侧主仓与昇腾侧一一对应，对照读就能理解移植差异。真正上机才能拿到的只剩性能手感与调参经验。
Q2： 有昇腾 950，第一个装哪个？
先装 TileKernels，它是唯一提供 pip 包的仓库，pip install tile-kernels 一条命令即可，再用仓库 pytest 验证。跑通后再源码安装 DeepGEMM-Ascend，最后才考虑 DeepEP-Ascend，因为后者还要求多卡、UBMEM 连通性与偶数 rank 条件。
Q3： clone 下来构建失败，最常见的原因是什么？
三个高频原因：子模块没拉全（DeepGEMM-Ascend 与 FlashMLA 要 recursive，DeepEP-Ascend 要单独拉 DeepJIT）；CANN 环境未激活或 ASCEND_HOME_PATH 没设；安装命令漏了 --no-build-isolation。按本文各路线的命令原文逐条核对，基本都能定位。
Q4： 昇腾侧的性能到底怎么样？
按 README 官方口径：FlashMLA 昇腾 prefill 达 410 TFlops（95% 硬件峰值）、decoding 达 360 TFlops（83% 硬件峰值）；DeepGEMM-Ascend 在多种矩阵形状上接近硬件极限性能。这些是官方自报数据，DeepEP-Ascend 的带宽测试条件 README 有完整描述，复现以仓库测试脚本为准。
Q5： DeepEP-Ascend 现在能直接用于生产吗？
建议缓一缓。一是仓库根目录截至 2026-10-01 快照没有 LICENSE 文件，商用合规待观察；二是验证栈边界很窄（950DT 加 CANN 9.2.0 那一套），偏离即无官方背书；三是推荐部署基线要等华为 Q3 商用 HDK 公开发布。学习与测试可以，生产引入等许可证与商用固件落地再评估。