1. 引言：为什么会有昇腾

2018 年前后，AI 算力需求开始以远超摩尔定律的速度增长。Transformer 架构在 2017 年被提出后，模型参数量从亿级迅速膨胀到千亿级，传统的通用 GPU 集群在能效比、互联带宽和供应链稳定性上都开始承压。与此同时，外部环境的变化让"算力自主可控"从一句口号变成了硬约束——高端 GPU 的供应不确定性，直接威胁到国内大模型训练与推理的连续性。

华为在这个背景下推出了昇腾（Ascend）系列 AI 处理器。它不是简单地对 GPU 做"国产替代"，而是从指令集、计算单元、互联拓扑到软件栈做了一套相对完整的自研体系。理解昇腾，不能只看单卡参数，而要把它当成"芯片 + 超节点 + 软件栈 + 编程范式"的四层结构来看。

本文按这个层次展开：先讲整体设计与产品谱系，再深入主力 910 系列的架构与参数，接着讨论实际训练/推理表现、超节点方案，最后落到 AscendC 编程范式、CANN 软件栈，以及对 PyTorch、Triton、TileLang 等上层框架的适配。

2. 昇腾的整体设计思路

昇腾系列的核心是华为自研的 达芬奇（Da Vinci）架构。与 GPU 的 SIMT（单指令多线程）模型不同，达芬奇采用 Cube + Vector + Scalar 三类计算单元 的异构组合：

• Cube 单元：专攻矩阵乘加（MAC），是深度学习中最耗算力的算子形态，用 3D Cube 结构在一个时钟周期内完成大规模矩阵运算。
• Vector 单元：负责向量运算，如激活函数、归一化、逐元素操作。
• Scalar 单元：负责标量控制流、地址计算和调度。

这种"专用单元做专用事"的设计，本质上是用更少的晶体管换取更高的矩阵运算能效比。代价是编程模型更复杂——你不能像写 CUDA 那样随意调度线程，而必须把算子拆解成 Cube/Vector 的流水。

在芯片之外，昇腾还定义了一套 Da Vinci 指令集，并围绕它构建了从驱动、编译器到算子库的完整软件栈。产品谱系大致分为两条线：

系列	定位	典型型号
昇腾 310	边缘/推理，低功耗	310、310P
昇腾 910	数据中心训练/推理	910A、910B（B1–B4）、910C
昇腾 950	新一代数据中心（2026 年上市）	950PR、950DT
昇腾 960	下一代（2027 年就绪）	960PR、960DT

从路线图看，华为已经把昇腾的演进节奏定为"一年一代"：910（2019）→ 950（2026）→ 960（2027）→ 970（2028 规划）。这个节奏背后是芯片、自研 HBM、超节点和 CANN 版本同步推进的系统工程，而不是单颗芯片的孤立迭代。

3. 主力 910 系列：架构与版本演进

910 是昇腾面向数据中心的主力产品，也是国内大模型训练集群中最常见的加速卡。它经历了多个版本的迭代，每一代在制程、算力、互联和显存上都有明显变化。

3.1 达芬奇核心的微架构

910 单芯片内部集成了多个 AI Core，每个 AI Core 又包含若干 Cube 和 Vector 单元。以 910B 为例，其单芯片集成了数十个 AI Core，通过片上网络（NoC）与 HBM 控制器、PCIe 控制器相连。关键设计点包括：

• 3D Cube 矩阵引擎：单周期可完成 16×16×16 的矩阵乘加，是 FP16/BF16 算力的主要来源。
• 多级片上缓存：L1/L0 缓存配合 Cube 的流水线，减少对 HBM 的访问压力。
• HBM 高带宽显存：910 系列普遍采用 HBM2/HBM2E，带宽达到 TB/s 级别，缓解大模型权重加载的瓶颈。

3.2 910A / 910B / 910C 的差异

三个版本的核心差异可以概括为"制程、算力、封装"三条线。需要先说明：910B 与 910C 的不少细节来自第三方分析或非官方渠道，不同信源口径存在差异，下表尽量标注区间。

版本	制程 / 封装	FP16 算力（单卡）	显存	互联 / 功耗
910A	台积电 7nm EUV	256 TFLOPS（INT8 512 TOPS）	32GB HBM2	HCCS / PCIe 4.0，最大 310W
910B	中芯国际 N+1（等效 7nm）	约 280–414 TFLOPS（随 B1–B4 型号而异）	64GB HBM2E	HCCS，带宽提升
910C	中芯国际 N+2，双 die 封装	约 800 TFLOPS（保守口径 640 TFLOPS）	96GB 或 128GB HBM2E，带宽约 1.8–3.2 TB/s	高速互联增强，约 310W

910A（2019） 是初代产品，基于台积电 7nm EUV 工艺和第一代达芬奇架构，单芯片集成 32 颗达芬奇 Max 核心，并集成了 CPU、DVPP（数字视觉预处理）等模块，支持 PCIe 4.0 与 HCCS 互联。它的历史意义在于：在 2019 年就把 FP16 算力做到了 256 TFLOPS 量级，同时把"训练卡"这个品类真正做成了可交付的产品。

910B 是在制裁背景下转向国内代工的一代，工艺换成中芯国际 N+1（等效 7nm），并按 B1/B2/B3/B4 细分出多个型号，算力与显存配置各不相同。架构上最重要的变化是 AIC（Cube 核）与 AIV（Vector 核）分离：Cube 核专职矩阵运算，Vector 核专职向量运算，两者通过片上总线协作，这让算子可以更精细地做流水编排，也让 AscendC 的"分离式编程"成为可能。

910C 走的是 Chiplet 双 die 封装 路线：把两颗 910B 级别的计算 die 封到一个封装内，晶体管规模约 530 亿，从而在单卡层面把算力堆到 800 TFLOPS 量级、显存堆到 96–128GB。这条路线的好处是不必等更先进制程，用封装换算力；代价是 die 间互联、功耗与散热设计难度显著上升，也让"单卡"这个概念在软件层面需要处理跨 die 的访存与通信。

需要说明的是，上表中的算力数字来自公开报道与行业分析，实际可用算力会因散热、功耗墙和软件调度而有所折扣。对 910B/910C 的精确核数、细分算力，不同信源尚未完全统一，工程选型时建议以华为官方或昇腾社区发布的技术文档为准。

3.3 910 的关键参数解读

对使用者而言，比"峰值 TFLOPS"更重要的参数是：

• 显存容量与带宽：决定能装下多大的模型、多长的上下文。64GB HBM2E 意味着单卡可以承载 7B 级模型的全参数推理，或作为大模型训练中的一个分片。
• 互联带宽：多卡训练时，AllReduce 等集合通信的瓶颈往往在卡间带宽。910 系列通过 HCCS（华为缓存一致性互联）提供远高于 PCIe 的卡间带宽。
• 功耗：单卡功耗在 300–400W 量级，直接决定了机柜的供电与散热设计。

4. 实际训练与推理表现

4.1 训练场景

在训练侧，昇腾的典型用法是构建 Atlas 训练集群，通过多机多卡并行训练大模型。实际表现受几个因素影响：

• 算子覆盖率：CANN 提供的算子库对主流模型（Transformer、CNN）覆盖较好，但遇到自定义算子时需要手写 AscendC，这会带来额外开发成本。
• 通信效率：在千卡规模下，集合通信的效率决定了线性加速比。昇腾通过 HCCL（华为集合通信库）做优化，实测在中等规模集群上可以接近线性扩展。
• 精度支持：FP16/BF16 是主力，部分场景支持 FP32 和 INT8。混合精度训练需要框架侧配合。

从公开的实践案例看，昇腾集群在训练百亿参数级模型时，单卡有效算力利用率（MFU）与同代 GPU 相比存在差距，但在特定优化过的模型上可以缩小到可接受范围。这个差距主要来自软件栈成熟度和算子调优深度，而非纯粹的硬件峰值。更具体地说，910B 时代的瓶颈集中在跨机通信：HCCS 在单机 8 卡内表现尚可，但跨机依赖 RoCE，AllReduce 长尾延迟可达 18ms 到 60–90ms 不等，长尾占比超过 25%，整集群同步被拖慢。这也是华为在 910C 时代转向超节点、把 384 卡收进一个高速互联域的直接原因（详见 5.3 节）。

4.2 推理场景

推理是昇腾相对更有优势的场景，原因在于：

• 推理对算子多样性的要求低于训练，主流模型的算子基本被覆盖。
• 昇腾 310/310P 在边缘推理上功耗优势明显。
• 910 系列配合 CANN 的图优化、量化工具，可以在 INT8 下获得较高的吞吐。

实际部署中，常见的做法是用 MindIE（昇腾推理引擎）做服务化，支持动态批处理、KV Cache 管理和多卡并行推理。

5. 超节点：把"卡"变成"一台机器"

单卡再强，也架不住大模型对显存和带宽的胃口。昇腾的方案是 超节点（SuperPoD）——通过高速互联把大量加速卡在逻辑上组成一台"超级计算机"。

5.1 架构要点

超节点的核心是 高速总线互联。以 910 时代的 Atlas 900 A3 SuperPoD（常被称为"昇腾 384 超节点"）为例，它的设计思路包括：

• 全互联拓扑：节点内的加速卡通过高速链路两两互联，总线带宽 1:1 无收敛，避免传统树形拓扑的带宽收敛。
• 统一内存寻址：通过互联协议让跨卡访问的延迟接近片内访问，从而支持模型并行时的频繁通信。
• 液冷散热：高密度集成带来巨大功耗，液冷成为标配。

需要强调的是，超节点不是"单颗芯片"，而是一个把多颗昇腾 910C（部分资料笼统标注为 910）通过高速总线互联、在逻辑上整合为一台"计算机"的大型系统。它由 12 个计算柜 + 4 个总线设备柜 组成，计算柜采用液冷散热，2025 年 3 月左右推出，满配支持 384 卡。华为云基于它构建的 CloudMatrix 384 云服务实例，在系统级算力上对标英伟达的 GB200 NVL72 系统。

5.2 参数与规模

超节点的关键指标是 节点内卡数、互联带宽、总显存。昇腾的超节点方案随芯片代际同步升级，规模跨度很大：

超节点	基于芯片	最大卡数	系统级算力	备注
Atlas 900 A3 SuperPoD	910C（部分资料标 910）	384 卡	307.2 / 288.7 PFLOPS @ FP16	2025 年 3 月推出，12 计算柜 + 4 总线柜，液冷
Atlas 950 SuperPoD	950DT	8192 卡（另有 1024 卡配置口径）	1 EFLOPS FP8 / 2 EFLOPS MXFP4	2025 年 9 月发布
昇腾 960 超节点	960	4096 卡	8 EFLOPS FP8	首个 NPO 近封装光学超节点

Atlas 900 A3 SuperPoD 的具体规格（华为官方与昇腾社区口径）：

• NPU 规模：最大支持 384 颗 昇腾 910 系列芯片；
• AI 算力：最大 307.2 / 288.7 PFLOPS @ FP16（具体数值因配置或统计口径略有差异）；
• 互联带宽：芯片间（D2D）双向互联带宽达 784 GB/s，总线带宽 1:1 无收敛；
• 内存能力：片上内存总容量达 48 TB（由 384 颗芯片各 128 GB 构成），支持统一编址，带宽速率最大 3.2 TB/s；
• 通信时延：单跳通信时延低至 200 ns。

关于 Atlas 950 SuperPoD 的卡数，公开信息存在两种口径：2025 年 9 月全联接大会发布时为 8192 卡，另有资料按 1024 卡 整机折算单卡算力（1024 卡、2 EFLOPS MXFP4，折合单颗 950DT 约 1.78 PFLOPS）。这可能是产品定义调整或不同配置版本所致，引用时建议注明来源与时间点。

超节点的意义不只是"卡多"。当节点内卡数达到数百乃至数千张、互联带宽达到数百 GB/s 到 TB/s 量级时，千亿甚至万亿参数模型的张量并行、专家并行可以尽量收敛在节点内完成，跨节点通信被压到更少的层数，这才是超节点相对普通集群的真正价值。

5.3 为什么必须造超节点：910 时代的通信瓶颈

理解了超节点的规格，还要理解它为什么会出现。在昇腾 910 时代（尤其是 910B），通信效率偏低确实是训练场景的一个核心短板，而且问题主要出在跨节点、跨机柜的互联上，而不是芯片本身的算力。

跨机通信是主要拖累。 910B 的片间高速总线 HCCS 在单机 8 卡内部表现尚可，双向带宽能与同期英伟达 A800 的 NVLink 持平。但一旦规模扩展到多机多卡，就得依赖 RoCE 网络。有开发者基于 910B 集群训练 70B 模型时反馈，AllReduce（梯度同步）的长尾延迟非常严重：同批次里部分卡只需 18ms，部分卡却要 60–90ms，长尾占比超过 25%，直接导致整集群同步阻塞，算力核心频繁空转。华为官方文档也记录过类似案例：因通信地址未对齐，HCCL 会额外插入一次内存拷贝，使单步通信时间从约 200ms 翻倍到 400ms。

互联拓扑存在先天限制。 910B 的 HCCS 采用对等拓扑，没有类似 NVSwitch 的专用交换芯片，机内 8 卡可以全互联，但跨机就得走外部网络。这导致大规模训练时，通信开销占比很容易被推高。华为自己的仿真数据也承认，传统 8 卡服务器堆叠的 10 万卡集群，集群内通信时间会超过训练总时间的 40%，严重制约算力利用率（MFU）。

华为的解法：用"超节点"换效率。 正因为看到了这个瓶颈，华为在 910C 时代主推的 Atlas 900 A3 SuperPoD（CloudMatrix 384 超节点），核心思路就是用更大的高速互联域来"吃掉"通信损耗：把 384 颗 910C 通过总线全对等互联，卡间带宽提升到 784 GB/s，单跳时延压到 200 ns。华为公布的数据是，相比传统 8 卡服务器集群，4K 超节点组成的 10 万卡集群 MFU 提升了 2.75 倍。

这相当于承认了旧架构的通信效率问题，并用系统级设计去弥补。所以更准确的表述是：昇腾 910 时代，单卡算力并不弱，但跨节点通信效率是训练大规模模型时的真实瓶颈，这也是华为后续从"堆服务器"转向"造超节点"的核心驱动力。

6. AscendC：为达芬奇而生的编程范式

CUDA 的成功很大程度上在于它给了开发者一个"够用"的编程模型。昇腾要建立生态，就必须回答"开发者怎么写算子"这个问题。答案是 AscendC。

AscendC 是华为推出的算子开发语言，本质上是 C++ 的扩展，通过内置的 API 和编程模型来调度 Cube/Vector 单元。它的核心概念包括：

• 核函数（Kernel）：运行在 AI Core 上的计算逻辑。
• 流水线（Pipeline）：显式地把数据搬运（CopyIn）、计算（Compute）、搬出（CopyOut）拆成多级流水，用异步方式掩盖访存延迟。
• 内存层级管理：开发者需要显式管理 Global Memory、L1、L0 之间的数据搬运。

一个简化的 AscendC 算子结构大致如下：

// 伪代码示意：AscendC 算子的典型三段式流水
extern "C" __global__ __aicore__ void add_custom(
    GM_ADDR x, GM_ADDR y, GM_ADDR z) {
    // 1. CopyIn：从全局内存搬到片上
    // 2. Compute：在 Vector/Cube 单元上计算
    // 3. CopyOut：结果写回全局内存
}

上面只是骨架。下面用一个真正能跑通的 向量加法（z = x + y） 算子，把三段式流水、Tiling 切分和内存层级都填进去。这个例子属于 Vector 类算子，不涉及 Cube，正好适合入门。

6.1 算子实现：Add 的三段式流水

#include "kernel_operator.h"

using namespace AscendC;

constexpr int32_t BUFFER_NUM = 2;   // 双缓冲，让搬运和计算重叠

class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}

    // 初始化：拿到全局内存地址、总长度、当前核负责的数据块
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z,
                                uint32_t totalLength, uint32_t tileLength) {
        // 每个 AI Core 只处理自己那一段，按 block 均分
        uint32_t blockIdx = GetBlockIdx();
        uint32_t offset = blockIdx * tileLength;

        xGm.SetGlobalBuffer((__gm__ float*)x + offset, tileLength);
        yGm.SetGlobalBuffer((__gm__ float*)y + offset, tileLength);
        zGm.SetGlobalBuffer((__gm__ float*)z + offset, tileLength);

        // 片上队列：每个 buffer 能放 tileLength 个 float
        pipe.InitBuffer(inQueueX, BUFFER_NUM, tileLength * sizeof(float));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, tileLength * sizeof(float));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, tileLength * sizeof(float));
    }

    __aicore__ inline void Process() {
        // 每个核处理 tileLength 个元素，按 32 个 float 一组循环
        int32_t loopCount = tileLength / 32;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    // 第一段：Global Memory -> Local Memory（片上）
    __aicore__ inline void CopyIn(int32_t progress) {
        LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
        LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
        // DataCopy 是异步搬运，配合队列实现流水
        DataCopy(xLocal, xGm[progress * 32], 32);
        DataCopy(yLocal, yGm[progress * 32], 32);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }

    // 第二段：在 Vector 单元上做逐元素加法
    __aicore__ inline void Compute(int32_t progress) {
        LocalTensor<float> xLocal = inQueueX.DeQue<float>();
        LocalTensor<float> yLocal = inQueueY.DeQue<float>();
        LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();

        Add(zLocal, xLocal, yLocal, 32);   // z = x + y，一次算 32 个

        outQueueZ.EnQue<float>(zLocal);
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }

    // 第三段：Local Memory -> Global Memory
    __aicore__ inline void CopyOut(int32_t progress) {
        LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
        DataCopy(zGm[progress * 32], zLocal, 32);
        outQueueZ.FreeTensor(zLocal);
    }

private:
    TPipe pipe;
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
    GlobalTensor<float> xGm, yGm, zGm;
};

// 核函数入口：每个 AI Core 启动一次
extern "C" __global__ __aicore__ void add_custom(
    GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tilingData, tiling);   // 从 Host 侧下发的切分参数
    KernelAdd op;
    op.Init(x, y, z, tilingData.totalLength, tilingData.tileLength);
    op.Process();
}

这段代码里有几个 AscendC 的关键点值得对照着看：

• GlobalTensor 与 LocalTensor：前者代表 HBM 上的全局内存，后者代表片上（L1/UB）的局部内存。算子必须显式地把数据从前者搬到后者，计算完再搬回去，这就是"显式内存层级管理"。
• TQue 队列 + BUFFER_NUM = 2：双缓冲让第 i 块的 Compute 和第 i+1 块的 CopyIn 并行，搬运延迟被计算掩盖，这是流水线的核心。
• EnQue / DeQue：入队出队不只是数据传递，还承担同步语义，保证搬运完成后再计算。
• GetBlockIdx()：每个 AI Core 拿到自己的编号，只处理属于自己的数据分片，天然支持多核并行。

6.2 Tiling：把大张量切成核能吃的块

上面的 tileLength 不是拍脑袋定的，它来自 Tiling。一个 1 亿元素的向量不可能一次塞进片上内存，Host 侧需要根据核数、片上 buffer 大小算出切分方案：

// Host 侧：计算 tiling 参数并下发
void ComputeTiling(uint32_t totalLength, uint32_t coreNum,
                   uint32_t& tileLength, uint32_t& tileNum) {
    // 每个核分到的元素数，向上取整
    tileLength = (totalLength + coreNum - 1) / coreNum;
    // 再按片上 buffer 容量切成更小的块
    tileNum = (tileLength + UB_BLOCK_SIZE - 1) / UB_BLOCK_SIZE;
}

Tiling 的本质是 在"核间并行"和"片上容量"之间做权衡：切得太粗，片上放不下；切得太细，循环开销和搬运次数上升。这也是 AscendC 相比 CUDA 更"重"的地方——CUDA 里线程块大小由硬件调度，AscendC 里切分策略要开发者自己算。

6.3 Host 侧调用：从框架到算子

算子写完后，还要通过 AscendCL 在 Host 侧完成内存分配、参数下发和 kernel 启动：

// Host 侧调用流程（简化）
aclrtMalloc(&xDevice, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
aclrtMalloc(&yDevice, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
aclrtMalloc(&zDevice, bytes, ACL_MEM_MALLOC_HUGE_FIRST);

aclrtMemcpy(xDevice, bytes, xHost, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
aclrtMemcpy(yDevice, bytes, yHost, bytes, ACL_MEMCPY_HOST_TO_DEVICE);

// 启动 kernel：blockDim 即使用的 AI Core 数
add_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice,
                                          workspace, tilingDevice);

aclrtSynchronizeStream(stream);
aclrtMemcpy(zHost, bytes, zDevice, bytes, ACL_MEMCPY_DEVICE_TO_HOST);

<<<blockDim, ...>>> 的写法刻意贴近 CUDA，降低迁移心智负担；但底层调度的是 AI Core 而非线程束，blockDim 对应实际启用的核数。

6.4 和 CUDA 写法对比

把同一个加法放到 CUDA 里，差异一目了然：

// CUDA：线程级并行，硬件自动管理缓存层级
__global__ void add_custom(const float* x, const float* y, float* z, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) z[i] = x[i] + y[i];
}

CUDA 版本里，开发者只描述"每个线程算什么"，缓存、搬运、同步都由硬件和运行时兜底；AscendC 版本里，搬运、缓冲、同步、切分全部显式暴露。这就是前文说的"学习曲线更陡，但能更精细控制硬件"的具体含义：多出来的代码量，换的是对流水和带宽的确定性掌控。

对于追求极致性能的算子，这种显式控制是必要的代价；而对于大多数业务算子，优先用 CANN 算子库或 Triton/TileLang 这类更高层的 DSL，只在热点算子上才下沉到 AscendC。

7. CANN：昇腾的软件栈底座

CANN（Compute Architecture for Neural Networks） 是昇腾的异构计算架构，地位类似 CUDA + cuDNN 的组合。它的分层大致是：

层次	组件	作用
应用层	MindSpore / PyTorch 适配	上层框架接入
图引擎	GE（Graph Engine）	图优化、算子融合、内存规划
算子层	AscendCL、算子库	算子调用与实现
运行时	Runtime	任务调度、流管理
驱动	Driver	硬件抽象

CANN 的关键能力包括：

• 图编译优化：把框架的计算图转成昇腾的图，做算子融合、常量折叠、内存复用。
• 算子库：提供覆盖主流模型的算子实现，减少手写 AscendC 的需求。
• 集合通信库 HCCL：多卡多机通信。
• 量化与调优工具：支持 INT8 量化、性能分析（Profiling）。

8. 框架适配：PyTorch 与 Triton / TileLang

8.1 PyTorch 适配

昇腾通过 torch_npu 插件接入 PyTorch。使用方式上，开发者只需把 .cuda() 换成 .npu()，大部分模型代码可以低改动迁移：

import torch
import torch_npu

device = "npu:0"
model = MyModel().to(device)
x = torch.randn(1, 3, 224, 224).to(device)
y = model(x)

torch_npu 负责把 PyTorch 的算子映射到 CANN 的算子，并处理设备管理、流同步等细节。实际迁移中，难点通常不在 API 替换，而在少数不支持的算子需要回退或手写实现。

8.2 Triton 与 TileLang 适配

Triton 是近年来在 GPU 生态中流行的算子 DSL，用 Python 写 kernel，由编译器生成底层代码。昇腾对 Triton 的适配思路是提供 后端编译器，把 Triton IR 降到昇腾的指令上。这能让熟悉 Triton 的开发者以较低成本在昇腾上写算子。

TileLang 则是更偏向"以 tile 为单位"的 DSL，强调用分块（tile）抽象来描述计算，天然契合 Cube 单元的矩阵分块运算。对昇腾而言，TileLang 这类 DSL 的价值在于：它把"如何切分矩阵、如何安排流水"这类硬件相关决策交给编译器，开发者只需描述 tile 级的计算逻辑。

这两条路线的共同目标是 降低算子开发门槛——AscendC 给了极致控制，Triton/TileLang 给了开发效率，形成梯度。

9. 昇腾 950：第三代达芬奇与双型号布局

昇腾 950 系列的信息在 2025 年 9 月 18 日的华为全联接大会 2025 上首次系统公布，轮值董事长徐直军披露了芯片演进路标与 950PR / 950DT 的规划及上市时间（分别为 2026 年 Q1 和 Q4）。与早期"参数不明"的印象不同，950 的架构与规格目前已经比较公开。

9.1 核心架构：第三代达芬奇 + Chiplet

昇腾 950 基于华为自研的 第三代达芬奇（Da Vinci）架构，并采用 多 Die 合封的 Chiplet 设计。以 950DT 为例，整颗芯片包含：

• 36 个 AI 子系统：每个子系统含 1 个 Cube Core 和 2 个 Vector Core，延续并强化了 910B 开始的 AIC/AIV 分离思路；
• 4 个 AI CPU Cluster：搭载自研 Linx816 CPU，负责控制流与调度；
• 128MB L2 Cache：单颗芯片集成，缓解大模型权重与 KV Cache 的访存压力；
• 精度支持：新增对 HiF8、MXFP8、MXFP4 等低精度格式的原生支持，直接面向大模型训练与推理的吞吐优化。

低精度原生支持是这一代最值得关注的变化之一。MXFP4 这类格式把单个参数的位宽压到 4 bit，配合分块缩放（block scaling）在精度可接受的前提下大幅提升吞吐，这正是当下大模型推理降本的主流方向。

9.2 950PR 与 950DT：同架构、不同内存

两款芯片共用同一套芯片架构，差异主要在搭载的 自研 HBM 上，从而面向不同场景：

项目	昇腾 950PR	昇腾 950DT
目标场景	推理 Prefill 阶段、推荐业务	推理 Decode 阶段、训练
自研内存	HiBL 1.0	HiZQ 2.0
容量	128GB	最高 144GB（部分配置 96GB）
带宽	1.6 TB/s	4 TB/s
算力峰值	MXFP4 约 1784 TFLOPS；FP8/HiF8 约 919 TFLOPS	MXFP4 约 1.78 PFLOPS（按 Atlas 950 整机折算）
互联	灵衢 2.0，双向最大 2 TB/s	灵衢 2.0，支持 PCIe 5.0、UBoE

这个分工很符合大模型推理的实际负载特征：Prefill 阶段 是一次性处理长 prompt，计算密集、对算力敏感，950PR 用较低带宽、较高算力的内存配置即可；Decode 阶段 是逐 token 生成，访存密集、对带宽和容量敏感，950DT 用 4 TB/s 的 HiZQ 2.0 把带宽拉满。训练场景同样吃带宽，因此也归到 950DT 一侧。

9.3 配套超节点

基于 950DT，华为推出了 Atlas 950 SuperPoD，最大支持 8192 张 NPU 高速互联（另有 1024 卡配置口径），提供 1 EFLOPS（FP8）或 2 EFLOPS（MXFP4）的系统级算力。它通过灵衢（UnifiedBus）网络把节点内互联延伸到集群层面，是 950 这一代"芯片—系统"协同的主要载体。

10. 昇腾 960 与后续路线

昇腾 960 的进展比原计划更快。根据华为在 2026 年 9 月全联接大会上公布的信息，昇腾 960 芯片性能实现倍增，研发进度超出预期：

• 昇腾 960DT：比原计划提前三个季度，预计 2027 年 Q1 就绪；
• 昇腾 960PR：比原计划提前一个季度，预计 2027 年 Q3 就绪。

在芯片就绪之前，基于 960 的 昇腾 960 超节点 已于 2026 年 9 月发布。它是首个采用 NPO（近封装光学） 技术的超节点，主打十万亿参数规模的大模型训练与推理，关键规格为：可扩展至 4096 卡，提供 8 EFLOPS FP8 算力，系统可用度达 99.8%；通过灵衢网络连接，最大可支持 100 万卡 的昇腾超节点集群。

NPO 的意义在于用光互联替代部分电互联，把长距离、高带宽、低功耗的互连能力推到封装附近，这是超节点从"千卡"走向"万卡、十万卡"的关键技术。按华为公布的节奏，后续的昇腾 970 计划在 2028 年推出，昇腾将坚持"一年一代"的演进。

11. 总结

昇腾的价值不能只用"单卡跑分"衡量。它是一套从达芬奇架构、910/950/960 系列芯片、超节点互联，到 AscendC 编程范式、CANN 软件栈、PyTorch/Triton/TileLang 适配的完整体系。这套体系目前的短板在软件生态成熟度和算子覆盖深度，优势在于自主可控、系统级协同和推理场景的能效比。

从代际演进看，昇腾的路线已经比较清晰：910 解决"有没有"，用 910C 的双 die 封装把算力堆上去，并用 Atlas 900 A3 SuperPoD 把 384 卡收进一个高速互联域来对冲跨机通信瓶颈；950 解决"好不好用"，用第三代达芬奇、自研 HBM 和 HiF8/MXFP4 低精度格式对齐大模型负载，并用 PR/DT 双型号区分 Prefill 与 Decode；960 则用 NPO 光互联把超节点推向万卡以上规模。
对开发者而言，务实的路径是：先用 torch_npu 跑通现有模型，遇到性能瓶颈再用 CANN 的 Profiling 定位，最后对热点算子用 AscendC 或 Triton 做定制。理解这套分层，比记住几个 TFLOPS 数字更有用。