📚 课程目录第2章 NPU基础

2、NPU基础：NPU架构特点、常见NPU类型（如华为Ascend、寒武纪、英伟达GPU对比）、NPU编程模型

NPU（Neural Processing Unit，神经网络处理器）是专为加速神经网络推理与训练而设计的专用处理器。与通用CPU和GPU不同，NPU在硬件层面直接针对矩阵乘加、卷积、激活函数、池化等深度学习核心算子进行优化，通过大规模并行计算阵列、片上高速缓存和专用数据通路，实现更高的算力密度与能效比。在TVM的NPU后端开发中，理解NPU的硬件架构、指令集特征和编程模型，是正确进行算子映射、内存规划和代码生成的前提。本章将系统梳理NPU的架构共性、主流NPU产品的设计差异，以及NPU编程模型的核心抽象，为后续TVM NPU后端的调度与代码生成打下基础。

2.1 NPU架构的核心设计特点

NPU的架构设计围绕神经网络计算的两个核心瓶颈展开：算力供给与数据搬运。与CPU追求低延迟、GPU追求通用并行不同，NPU在硬件层面做了大量针对性的“减法”与“加法”。理解这些特点，是后续在TVM中编写NPU后端调度策略的基础。

2.1.1 核心原理：从指令流到数据流的转变

传统CPU采用冯·诺依曼架构，指令串行执行，算力受限于指令发射宽度。GPU通过SIMT（单指令多线程）模式提升了并行度，但依然保留了复杂的线程调度与缓存一致性开销。NPU则普遍采用数据流驱动或脉动阵列架构，将计算映射为数据在硬件流水线中的流动。以华为Ascend的达芬奇架构为例，其AI Core包含三个基础单元：Cube单元（负责矩阵乘加）、Vector单元（负责向量运算）、Scalar单元（负责标量控制）。这种异构计算单元的组合，使得NPU在执行卷积时，能够将输入特征图直接“灌入”Cube阵列，通过硬件级流水线完成乘加累加，而无需像GPU那样频繁访问寄存器文件。

从TVM的角度看，这种架构意味着算子调度必须显式管理数据在SRAM与外部存储之间的搬移。TVM的调度原语（如copy、compute_at）需要映射到NPU的DMA引擎和片上缓存管理指令上。

2.1.2 判断方法：识别NPU架构的关键指标

在评估一款NPU或为其编写TVM后端时，应关注以下架构特征：

MAC阵列规模与组织方式：例如16x16、32x32的脉动阵列，决定了单周期矩阵运算的维度。
片上存储层次：包括L0 Buffer、L1 Cache、共享内存（Shared Memory）的容量与带宽。TVM的storage_align和pragma需要针对这些层次进行优化。
数据通路宽度：DMA搬运的位宽（如128bit、256bit）直接影响带宽利用率。
指令集抽象层级：是暴露底层硬件细节（如Ascend的TBE指令），还是提供高层算子库（如寒武纪的NeuWare）。

经验提醒：在TVM中编写NPU后端时，不要试图用GPU的线程模型去理解NPU。NPU通常采用“任务级并行”而非“线程级并行”。TVM的thread_binding在NPU后端往往映射为“核间并行”或“流水线级并行”，而非GPU的block/thread概念。

2.1.3 案例拆解：Ascend 310 的AI Core流水线

以华为Ascend 310为例，其AI Core执行一个卷积算子时，TVM生成的指令序列大致如下：

// 伪代码示意：TVM调度映射到Ascend指令
// 1. 数据搬运：将输入从DDR搬运到L1 Buffer
dma_copy(dst=L1_Buffer, src=DDR_Addr, size=Input_Tensor_Size);
// 2. 矩阵计算：Cube单元执行矩阵乘加
cube_instruction(operand_a=L1_Buffer, operand_b=Weight_Buffer, output=L0_Buffer);
// 3. 向量运算：激活函数或池化
vector_instruction(op=RELU, input=L0_Buffer, output=L0_Buffer);
// 4. 数据写回：将结果从L0写回DDR
dma_copy(dst=DDR_Out_Addr, src=L0_Buffer, size=Output_Tensor_Size);

这个流程体现了NPU架构的典型特征：计算与搬运分离。TVM的调度器需要确保在Cube单元计算当前数据块时，DMA引擎已经在搬运下一块数据，从而实现流水线重叠。如果TVM后端未能正确生成异步DMA指令，NPU的算力利用率可能不足30%。

2.1.4 常见误区与总结

风险提示：初学者常误以为NPU的“算子库”是万能的。实际上，TVM的价值在于绕过固定算子库，为自定义算子生成高效代码。但NPU硬件通常对算子形状有严格约束（如对齐要求）。若TVM生成的指令未满足对齐条件（如通道数不是16的倍数），会导致硬件报错或性能骤降。

总结：NPU架构的核心是“算力专用化”与“数据流优化”。TVM NPU后端的首要任务，是将高层IR映射到这种数据流执行模型上。

2.2 常见NPU类型对比：华为Ascend、寒武纪、英伟达GPU

虽然都用于加速神经网络，但不同NPU的设计哲学差异巨大。下表从TVM后端开发最关心的维度进行对比：

对比维度	华为 Ascend (达芬奇)	寒武纪 MLU (思元)	英伟达 GPU (Tensor Core)
计算核心	AI Core（Cube+Vector+Scalar）	MLU Core（矩阵+向量单元）	SM（流式多处理器）+ Tensor Core
编程模型	TBE（Tensor Boost Engine）算子开发，DSL	NeuWare SDK，支持C++/Python	CUDA C/C++，PTX，cuDNN
TVM支持	TVM原生支持（通过BYOC或Codegen）	TVM社区有贡献，但成熟度较低	TVM支持最完善（CUDA后端）
内存模型	统一寻址，但需显式管理L1/L0 Buffer	片上SRAM + 全局内存，需手动管理	全局内存 + 共享内存 + 寄存器，层次清晰
指令粒度	粗粒度（算子级指令）	中粒度（指令级并行）	细粒度（线程级指令）
典型应用	端侧推理、边缘计算	云端训练、推理	通用计算、训练、推理

2.2.1 核心原理：三种架构的算力组织差异

英伟达GPU本质上是通用并行处理器，Tensor Core只是其SIMT架构上的一个“加速附件”。因此，GPU的编程模型依然以线程为中心，TVM的CUDA后端可以沿用blockIdx、threadIdx等概念。而Ascend和寒武纪则是专用数据流处理器，其编程模型更接近于“描述数据如何流动”而非“描述线程如何执行”。例如，Ascend的TBE编程中，开发者需要明确指定输入张量放在哪个Buffer，计算单元如何读取。

2.2.2 判断方法：如何选择TVM后端策略

在TVM中为不同NPU编写后端时，应遵循以下判断路径：

检查硬件是否支持TVM的标准算子库：若支持（如GPU），可直接复用tvm.contrib.cudnn。
若为专用NPU：优先考虑TVM的BYOC（Bring Your Own Codegen）框架，将子图下沉到厂商SDK。
评估算子粒度：Ascend适合粗粒度算子融合，GPU适合细粒度线程调度。

2.2.3 案例拆解：同一卷积在三种硬件上的TVM调度差异

假设我们要在TVM中调度一个3x3卷积，输入为1x64x112x112：

GPU (CUDA)：TVM会生成threadIdx.x绑定到输出通道，blockIdx.x绑定到输出像素块，利用共享内存做数据复用。
Ascend：TVM会生成dma_copy指令将输入分块搬运到L1，然后调用cube_instruction执行矩阵乘，最后用vector_instruction做bias add。
寒武纪：TVM会尝试调用NeuWare的cnml_convolution_forward算子，若算子形状不匹配，则回退到逐元素指令。

这个对比说明，TVM NPU后端的核心工作不是“生成代码”，而是“映射到正确的硬件抽象层”。

2.2.4 常见误区与总结

复盘要点：不要将“NPU”视为一个统一的硬件类别。Ascend的达芬奇架构与寒武纪的MLU架构在指令集、内存模型上几乎没有共性。TVM的抽象层（Relay/TE）必须通过不同的Codegen和Runtime来适配。在课程后续章节中，我们将分别实现针对Ascend的codegen_ascend.cc和针对寒武纪的codegen_cambricon.cc。

2.3 NPU编程模型：从Host到Device的抽象

NPU编程模型定义了主机（Host）侧如何描述计算任务、如何管理设备（Device）侧的内存与执行流。TVM的NPU后端必须遵循这些模型，才能正确生成可执行的二进制或运行时调用。

2.3.1 核心原理：任务描述与内存管理

所有NPU编程模型都包含三个核心抽象：张量（Tensor）描述符、算子（Operator）指令、执行流（Stream）。以寒武纪为例，其编程模型要求：

// 寒武纪 NeuWare 伪代码
cnmlCreateTensorDesc(&input_desc, CNML_DTYPE_FLOAT16, dims, 4);
cnmlCreateConvOp(&conv_op, input_desc, weight_desc, conv_param);
cnmlSetOpWorkingSpace(conv_op, workspace_size);
cnmlComputeConvOpForward(conv_op, input_ptr, weight_ptr, output_ptr, stream);

TVM在对接这类SDK时，需要将Relay的nn.conv2d节点转换为上述API调用序列。关键在于内存布局转换：NPU通常要求NHWC或特殊的5HD格式（Ascend），而TVM内部默认NCHW。因此，TVM后端必须插入layout_transform算子。

2.3.2 操作流程：TVM NPU后端的标准对接流程

算子识别：在Relay中遍历图，识别可下沉到NPU的算子子图。
布局转换：将NCHW转换为NPU原生布局。
内存规划：调用厂商SDK的内存分配接口，或直接操作设备指针。
指令生成：将TVM的TE调度转换为厂商指令或SDK调用。
运行时封装：实现TVM的PackedFunc，将NPU执行封装为TVM可调用的函数。

2.3.3 案例拆解：Ascend的TBE与TVM的融合

华为Ascend提供了TBE（Tensor Boost Engine）作为算子开发工具。TBE支持使用Python编写算子，并通过DSL描述计算逻辑。TVM的Ascend后端有两种接入方式：

方式一（推荐）：使用TVM的BYOC框架，将整个子图编译为TBE算子。
方式二：使用TVM的TE调度，通过tvm.contrib.ascend接口逐指令生成。

方式一更稳定，但灵活性差；方式二灵活，但需要深入理解Ascend的指令语义。在实际项目中，通常采用混合策略：标准算子走BYOC，自定义算子走TE。

2.3.4 常见误区与总结

风险提示：NPU编程模型中的“同步”与“异步”概念极易混淆。例如，Ascend的dma_copy是异步的，必须在计算指令前插入sync指令。TVM的调度器若未正确插入同步屏障，会导致数据竞争。建议在TVM后端中，将同步逻辑封装为独立的tir.call_extern调用。

总结：NPU编程模型的核心是“显式数据流管理”。TVM NPU后端的开发，本质上是将TVM的调度抽象（如compute_at、storage_align）翻译为NPU的“搬运+计算+同步”指令序列。掌握本章内容后，读者应能理解为何TVM的NPU后端不能简单复用CUDA代码生成逻辑，而必须针对每个硬件平台定制化开发。