CANN Ascend C语言扩展深度解读：SIMD/SIMT混合编程模型与Reg向量化架构设计原理

前言

做深度学习框架开发的人，绕不开一个问题：芯片厂商提供的算子开发工具，到底是一套什么样的接口？在 NVIDIA 的生态里，这个答案是 CUDA——从硬件抽象到编程模型到工具链，一整套完整体系。昇腾 CANN 也有自己的答案，这套答案的名字叫 Ascend C。

asc-devkit 是昇腾 CANN 推出的官方算子开发语言仓库，定位是"昇腾NPU专用的算子程序开发语言"，v9.0 新增的 SIMD/SIMT 混合编程和 Reg 寄存器级编程能力，是昇腾 950PR 硬件特性的第一波软件表达。v9.0 是一个值得注意的节点版本，因为它把 SIMD/SIMT 混合编程能力和 Reg 寄存器级编程这两件事第一次摆到了开发者的面前。这两件事分开来看各有价值，放在一起看，实际上是在回答一个根本性的问题：当昇腾 NPU 的硬件能力越来越丰富之后，开发者应该用什么样的抽象来驾驭它？

这篇文章从 Ascend C 的设计哲学出发，先说清楚"类库加语言扩展"这两层架构是怎么分工的，然后用 Reg 编程模型作为切入点解释寄存器级操作的意义，再来看 SIMT 怎么补上了 warp 级并行的短板，最后落到 SIMD/SIMT 混合编程的具体场景和 MXFP 低比特矩阵乘的新接口。目标不是堆砌 API 清单，而是把每一层抽象出现的动机说清楚，让读者理解之后能自己判断在什么场景该选什么接口。

Ascend C 的定位再认知：从 CUDA/HIP 出发理解"类库+语言扩展"的两层设计

在说 Ascend C 之前，有必要先建立一点横向的坐标感。CUDA 从诞生之初就走的是编译器加运行时加类库三合一的路子，开发者写的 kernel 代码最终编译成 PTX 再映射到真实硬件。HIP 则是一个接口兼容层，把 CUDA 代码翻译成能在 AMD GPU 上跑的 HIP 代码。两者在编程模型上是一套，但在硬件抽象层走的路径不同。

昇腾 CANN 的做法跟这两者都有区别。CANN 本身是昇腾异构计算架构，包含编译器、运行时、算子库、通信库等一整套软件栈。Ascend C 位于这个架构的第二层，也就是"昇腾计算服务层"，具体落在"算子开发接口"这个位置。它不是 CANN 的全部，而是 CANN 里面专门负责"怎么写一个算子"的那一层。

这种定位决定了 Ascend C 的设计哲学不是"我要做一套全新的编程语言"，而是"我要在 C 和 C++ 标准规范的基础上，提供一套够到底层的接口"。原生支持 C 和 C++ 标准规范这句话不是宣传语，它意味着开发者可以用自己熟悉的指针操作和数组语法来写昇腾 NPU 上的代码，不需要学一套全新的方言。类库层和语言扩展层的两层设计，也是围绕这个思路展开的。

类库层（基础 API 加高阶 API）是一套 C++ 类库，基于 Tensor 编程，提供单指令抽象。开发者用 MakeTensor 分配内存，用 LocalMemoryAllocator 管理显存，调用封装好的算子接口。这个层级的优点是语义清晰、上手快，适合需要快速验证算法的场景。语言扩展层则是 SIMD 和 SIMT 的 API，用纯 C 接口提供，开发者用数组语法分配内存，直接操作指针。这两个层级不是割裂的，而是给不同技能栈和性能目标的开发者提供了不同的入口。

理解这个分层之后，选择接口就不再是盲目的。如果你熟悉 Tensor 编程、想快速验证想法，从基础 API 入手。如果你在做极致性能优化、需要精确控制每一次内存搬运和每一条指令的调度，语言扩展层的 SIMD/SIMT API 才能满足要求。两者之间没有绝对的优劣，只有场景的匹配程度不同。

Ascend C 的设计目标里有五个关键词：高性能、完备性、易编程、可调试、兼容性。高性能对应的是开放芯片的完备编程能力，让专家级开发者能挖掘硬件的全部潜力。完备性是说三层 API（TPipe/TQue 框架编程 API、基础 API、语言扩展层 SIMD&SIMT API）都能实现底层的完备编程能力，不存在"这个操作只能在更高层做"的能力天花板。易编程和可调试则通过框架层自动管理同步与内存、通过 printf 和 DumpTensor 等工具来保障。兼容性体现在对 C/C++ 标准的最小化扩展，以及 Atlas A2/A3/950PR/950DT 多代硬件之间的 API 跨代兼容能力上。

下面给出一个 SIMD 基础 API 调用的示例，展示向量加法这种最朴素场景下 Ascend C 的写法。这个例子来自 asc-devkit 仓库中03_basic_api目录下的样例逻辑，展示了基础 API 中 Vector 计算接口的典型调用方式。

#include"operator/operator.h"#include"operator/variable.h"classVecAddKernel:publicKernelBuilder{public:__aicore__inlineVecAddKernel(){}__aicore__inlinevoidInit(KernelInput&input,KernelOutput&output,KernelParam&param,KernelContext&ctx){a.Init(input.a_global,param.aLength);b.Init(input.b_global,param.bLength);out.Init(output.y_global,param.outLength);}__aicore__inlinevoidProcess(){// 每次处理一个Tile的数据constexpruint32_tBLOCK_LENGTH=256;uint32_ttileNum=(param.aLength+BLOCK_LENGTH-1)/BLOCK_LENGTH;for(uint32_ti=0;i<tileNum;++i){// 从 GlobalMemory 搬到 Unified Buffera.SetGlobalBuffer(input.a_global,BLOCK_LENGTH);b.SetGlobalBuffer(input.b_global,BLOCK_LENGTH);// 调用基础API中的Vector加法接口// 这里用的是Vec::Add，底层映射到SIMD向量指令Vec::Add(out,a,b,BLOCK_LENGTH);// 同步一条，确保当前Tile完成再处理下一个Wait();}}private:Tensor<>a,b,out;KernelContext ctx;};

为什么不直接 for 循环逐元素加，而要用 Vec::Add？因为 Vec::Add 底层对应的是 Vector Core 的 SIMD 指令，一次性对 128 位甚至更宽的寄存器做多个元素的并行计算。SetGlobalBuffer 这类接口做的事情，是把数据从 GlobalMemory 通过 DMA 搬到 Unified Buffer 的特定位置，Unified Buffer 在 Vector Core 内部，是离计算单元最近的高速存储区域。把数据先归集到 Unified Buffer 再做计算，而不是直接跨地址随机访问，是昇腾 AI Core 的 Vector 编程范式里最核心的一条规则。Wait() 同步在每个 Tile 边界出现，是因为 DMA 搬运和 Vector 计算在硬件上是两条并行的流水线，如果不加同步就接着处理下一个 Tile，可能会出现数据相关性导致的错误。

Reg编程模型：绕过Unified Buffer的寄存器级编程

在说 Reg 编程之前，先把 Reg 这个名字拆开理解一下。Reg 不是"寄存器优化"那种营销词汇，它指的是 RegBase 架构——昇腾 950PR 上新增的一种底层执行模式，对应的是 Vector 寄存器直接寻址的编程路径。

在传统的基础 API 模型里，数据必须经过 GlobalMemory 到 Unified Buffer 再到 Vector 寄存器的路径，计算结果写回 Unified Buffer 再写回 GlobalMemory。Unified Buffer 是一个共享的高速存储区域，所有的计算单元共享这块空间，硬件会自动管理缓存一致性。这套机制在大多数场景下工作得很好，但它有一个隐含的约束：当你需要做细粒度的寄存器复用、或者需要绕过缓存一致性开销做极致优化的时候，Unified Buffer 的间接层就成了负担。

Reg 编程模型的核心变化，是提供了一套 API 让开发者直接操作 Vector 寄存器，不再需要经过 Unified Buffer 这一层中介。根据 v9.0 的文档，Reg 编程接口目前提供 90 多个，涵盖 Reg 数据搬运、基础算术运算、规约计算、同步控制等类别。这些接口位于impl/basic_api/reg_compute/dav_3510目录下，从文件名就能看出它们是面向 Ascend 950PR（也就是 dav_3510 架构）的。

寄存器级编程跟统一缓存编程的本质区别，可以通过一个生活中的类比来理解。Unified Buffer 像是一个公共厨房，所有厨师都在这个厨房里拿食材、做菜、还食材。好处是不用自己买冰箱，坏处是高峰期大家要排队，而且公共厨房的储物格大小是固定的。Reg 编程则是每个厨师带着自己的随身保温箱，自己管自己的食材，想拿什么直接拿，不用跟别人打招呼。

这种模式适合什么场景呢？适合那些数据局部性极强、一次只处理很少数据但需要反复复用这些数据的计算任务。典型的场景包括激活函数里对单个元素的非线性变换、规约计算里的分块求和、以及需要精确控制指令流水线的访存密集型算子。

下面是一个 Reg 编程接口的示例结构，展示了如何用寄存器级接口做矢量计算。这个示例逻辑来自仓库中 reg_compute 目录下的实现文件，展示了 VecBinary、VecUnary 等典型 Reg 接口的调用模式。

#include"kernel_operator.h"// Reg编程接口直接操作Vector寄存器，不需要通过Unified Buffer// 这里展示的是注册级矢量二元运算的典型写法classRegVecAddKernel{public:__aicore__inlineRegVecAddKernel(){}__aicore__inlinevoidInit(RegAddr&addr,uint32_tlen){this->len=len;// Reg编程模式下，通过RegAddr初始化寄存器地址// 不需要像基础API那样先做SetGlobalBufferregA.Init(addr.GetSrc0Addr(),addr.GetSrc0Size());regB.Init(addr.GetSrc1Addr(),addr.GetSrc1Size());regC.Init(addr.GetDstAddr(),addr.GetDstSize());}__aicore__inlinevoidProcess(){// VecBinary::Add 是寄存器级矢量加法// 底层直接发射 SIMD 指令到 Vector Core，不经过 Unified Buffer// 这里的 Register 参数是直接映射到硬件寄存器的VecBinary::Add(regC,regA,regB,len);// Reg编程不需要显式 Wait()，因为没有 DMA 流水线需要同步// 但如果涉及跨 Tile 操作，仍需要 Membar 同步指令if(len>THRESHOLD){Membar::Sync();}}private:uint32_tlen;// Reg 编程中使用 VecBinary、VecUnary 等命名空间下的接口// 这些接口直接接收寄存器地址作为参数VecBinary::RegT regA,regB,regC;};

寄存器级编程省掉了 Unified Buffer 这一层，代价是开发者要自己承担更多的内存管理职责。在基础 API 里，硬件会自动把数据从 GlobalMemory 预取到 Unified Buffer 并维护一致性，开发者不需要关心这块数据是在寄存器里还是缓存里。Reg 编程把这层透明性打破了，开发者要直接知道"我的数据现在在哪个寄存器"、“这些寄存器够不够用”。这种设计的回报是零拷贝开销和更高的寄存器复用率——对于那些单次计算量很小但调用频次极高的算子，Reg 编程能显著减少数据在存储层次之间的搬运次数。Membar::Sync() 在 len 超过某个阈值时才调用，是因为在极短长度下，同步指令的延迟可能反而超过它能解决的问题。

Reg 编程和传统 TPipe/TQue 框架编程模型的关系，也值得在这里辨析一下。TPipe/TQue 是 Ascend C 早期版本就有的高层编程框架，TPipe 管理计算流水线，TQue 提供队列式的任务调度，两者配合可以实现复杂的算子融合和流水线并行。这套框架的优势是开发者不需要关心同步细节，框架会自动插入必要的 barrier。Reg 编程则是彻底绕开这套框架，直接用底层的寄存器操作接口。从 TPipe/TQue 到 Reg，是从"让框架帮你管同步"到"你自己对每一纳秒负责"的跨越。

SIMT与SIMD混合编程：Atlas 950PR的warp级并行与SIMD向量指令协同

SIMT 的全称是 Single Instruction Multiple Thread，翻译过来是单指令多线程。这是 NVIDIA GPU 编程模型里的核心概念，但在昇腾 NPU 上引入 SIMT 是 v9.0 的新动作。

理解 SIMT 的意义，先要从 SIMD 说起。SIMD（Single Instruction Multiple Data）是单指令多数据的意思，一条指令同时对多个数据元素进行操作。比如一条 256 位的向量加法指令，可以同时处理 8 个 32 位浮点数。SIMD 的优点是效率高，缺点是每条线程执行的指令序列完全相同，无法处理分支。遇到 if-else 这种分支的时候，SIMD 的处理方式是执行两条路径然后用 mask 把不需要的结果置零，这在发散控制流场景下会产生巨大的浪费。

SIMT 对这个问题的解法是引入线程级并行。每条线程有自己的 PC（程序计数器），可以独立寻址、独立走不同的执行路径。硬件上，多条线程被编成 warp（SIMT 中的线程束概念），warp 内的线程以 SIMD 方式同步执行相同的指令，但当遇到分支时，硬件会自动屏蔽不需要执行当前路径的线程。这样做的好处是，即使代码里有分支，每条线程也只需要执行自己那条路径上的指令，而不用像纯 SIMD 那样两条路都走一遍。

在昇腾 950PR 的 SIMT 编程模型里，硬件抽象出的并行层次是这样的：Vector Core 内部有多个线程块（thread block），每个线程块包含多条线程，线程块内所有线程共享一块 Unified Buffer 空间作为 Shared Memory。Global Memory 是核外全局内存，被所有 Vector Core 共享。L2 Cache 是所有 Vector Core 共享的高速缓存，位于 Global Memory 与 Data Cache 之间，由硬件自动管理。Data Cache 则是 Unified Buffer 中预留出来做数据中转的空间，最大 128KB 最小 32KB，大小由开发者自主分配。

SIMT API 分布在include/simt_api目录下，包含约 700 个接口，按功能分为 warp 级别、原子操作、基本数学计算、类型转换等类别。关键的 API 头文件包括device_warp_functions.h（warp 级操作）、device_atomic_functions.h（原子操作）、device_functions.h（通用设备函数）、vector_functions.h（向量函数）、cooperative_groups.h（协作组）等。

在 Atlas 950PR 上，SIMD 和 SIMT 不是非此即彼的关系，而是可以混合使用。SIMD 适合规则的大块数据并行，SIMT 适合不规则的、带有分支的并行计算。混合编程的场景通常是这样：整体任务划分用 SIMT 线程块来做，每个线程块内部再细分为 SIMD 向量操作。打个比方，SIMT 像是一群工人分组去搬运不同的货物，SIMD 像是每个工人同时用多只手抓起多个箱子。两者协同的前提，是理解各自擅长什么——SIMT 管"谁做什么任务"，SIMD 管"每个任务怎么做"。

下面给出一个 SIMT 编程的基础示例，展示了 SIMT 模式下如何用线程级并行实现 Gather 操作。这个示例逻辑来自 asc-devkit 中 SIMT 样例目录下的实现。

#include"simt_api/asc_simt.h"usingnamespaceAscendC;// SIMT kernel：每个线程独立计算自己的输出extern"C"__global__ __SIMT__voidgather_1d_custom(float*input,int32_t*index,float*output,int32_tout_size){// 每个线程根据 threadIdx.x 和 blockIdx.x 计算自己的全局线程IDuint32_tgid=blockIdx.x*blockDim.x+threadIdx.x;if(gid>=out_size){return;}// SIMT 模式下，每条线程可以独立访问任意全局内存地址// 这是 SIMT 和 SIMD 最大的区别：SIMD 必须是规则的数据布局// SIMT 允许线程访问任意索引的数据，不要求连续int32_tidx=index[gid];floatval=input[idx];// 写入对应的输出位置output[gid]=val;}// Host端调用，使用标准的 CUDA/HIP 风格语法// blocks_per_grid 和 threads_per_block 由开发者根据数据规模配置voidlaunch_gather(std::vector<float>&input,std::vector<int32_t>&index,std::vector<float>&output,void*stream){uint32_tblocks=(output.size()+255)/256;uint32_tthreads=256;// <<<gridDim, blockDim, sharedMem, stream>>> 是 SIMT 编程的标准调用方式gather_1d_custom<<<blocks,threads,0,stream>>>(input.data(),index.data(),output.data(),output.size());}

Gather 操作天然是不规则访问——每个线程要取的数据位置由 index 数组决定，各线程之间没有连续地址的约束。用 SIMD 来做这件事，需要先把所有要取的数据地址算出来，然后用 gather 指令一条一条取，效率很差。SIMT 的每个线程有独立的 PC，可以独立算自己的 index 然后直接访问，线程之间互不干扰。__SIMT__标记说明这个 kernel 运行在 SIMT 模式下，硬件会为每个线程块分配独立的寄存器组，线程块内的所有线程共享一块 Shared Memory（位于 Unified Buffer）。Blocks 和 threads 的配置是有讲究的：threads_per_block 越多，每个线程能分到的寄存器数量越少；threads_per_block 越少，线程块间的调度开销占比越大。

下面给出一个效率对比表，展示在 Atlas 950PR 上分别使用纯 SIMD、纯 SIMT、以及 SIMD/SIMT 混合编程三种模式时，在不同类型算子上的行为差异。注意：这里的数据是概括性描述，具体的数值会因硬件配置和算子实现而有差异。

MXFP低比特矩阵乘扩展：高阶Matmul API对MXFP4/8的支持

矩阵乘是深度学习里最核心的计算内核，几乎所有的神经网络层最终都可以拆解为矩阵乘或其变体。在昇腾 NPU 上，矩阵乘一直是高阶 API 中的重点能力，v9.0 里的新动作是把 MXFP4 和 MXFP8 这两种低比特数据类型的矩阵运算做进了 Matmul 高阶 API。

MXFP 是混合精度浮点格式的简称。FP4 和 FP8 是极低比特宽度的数值表示，用 4 位或 8 位来表示一个浮点数。问题在于，4 位浮点无法像 FP16/BF16 那样用标准 IEEE 754 格式来定义，因为 4 位只能表示 16 个不同的值。MXFP 的做法是混合精度：输入和权重用低比特表示（MXFP4/MXFP8），累加过程用高精度（通常是 FP32），输出再量化为低比特。这样既能把矩阵乘法的数据量压缩到原来的 1/4（MXFP4）或 1/2（MXFP8），又不至于在累加过程中因为精度损失而导致模型效果明显下降。

MXFP4/8 对应的实际价值有两个层面。第一层是内存占用减半或更多，对于显存敏感的推理场景来说意义明显。第二层是算力吞吐倍增——当数据类型变窄之后，同一块芯片在单位时间内能处理的矩阵元素数量翻倍，对应的就是吞吐量提升。当然，这两个收益的前提是算法对低精度不敏感，这通常是经过量化感知训练之后才能保证的。

在 asc-devkit 的impl/adv_api/detail/matmul/mx_matmul_impl.h文件里，可以看到 MXFP Matmul 的实现模板。这个实现文件不是一个可直接调用的接口，而是一个模板实现，真正的公开接口是adv_api/matmul/matmul.h中定义的Matmul类模板。从实现代码里可以看到 MXFP Matmul 的几个关键步骤：CopyCubeInScaleA 和 CopyCubeInScaleB 做的是权重的缩放和打包，把 MXFP4/MXFP8 格式的数据转换成内部可用的表示；MatmulTensorInfoScaleA 和 MatmulTensorInfoScaleB 管理的是缩放因子和原始数据之间的对应关系；实际的矩阵乘核心由底层的 Cube 单元完成。

对普通开发者来说，理解 MXFP Matmul 的意义不在于去读懂模板实现里的每一个细节，而在于知道"我可以在 Ascend C 的高阶 API 里直接用 MXFP4/MXFP8 来跑矩阵乘"这件事是真的，并且知道这个能力是怎么被集成进来的。当你在用 ATB（Ascend Transformer Boost）或者其他基于 Ascend C 的加速库时，如果底层调用的 Matmul 接口支持 MXFP4/8，你就自动享受到了这套能力，不需要自己手写量化逻辑。

结尾

回顾这篇文章的脉络，其实就是在回答一个问题：Ascend C 的三层 API 加两套执行模式，是怎么组织在一起的，以及什么时候该用哪一层。

最顶层的 TPipe/TQue 框架 API 帮你包办了同步和内存管理，你只需要写计算逻辑，适合想快速出结果、不想纠缠底层细节的时候。基础 API 把同步和内存的掌控权还给你，但接口语义还是围绕 Tensor 展开的，适合对性能有要求但不需要精确到每条指令的场景。语言扩展层的 SIMD API 则更进一步，用纯 C 接口和数组语法提供最直接的硬件编程能力，适合需要在向量计算上做极致优化的专家。Reg 编程模型是在 SIMD 基础上又往前走了一步，把 Unified Buffer 也绕过去，直接操作 Vector 寄存器。SIMT 编程模型则是另一条路，引入线程级并行来处理 SIMD 不擅长的不规则访问和分支发散场景。

SIMD/SIMT 混合编程实际上是昇腾 950PR 上最强大的编程范式——不是选 SIMD 或者选 SIMT，而是识别出一个算子里哪些部分是规则的批量操作、哪些部分是离散的索引操作，分别用最适合的工具来处理。MXFP 低比特矩阵乘则是这个能力在 Transformer 时代的一个具体应用，数据格式变了，但调用的接口还是同一套 Matmul 高阶 API。

昇腾 NPU 从 Ascend 910 到 Ascend 950PR 的演进，硬件能力在不断丰富。CANN 在这个演进过程中的角色，是确保每一代硬件的能力都能被相应的 API 层级所表达出来，既不让专家级开发者感到束手束脚，也不让普通开发者被复杂性吓退。asc-devkit 仓库里的 260 个样例，就是这套 API 体系的活教材，从 SIMD 向量加法到 SIMT Gather 再到 MXFP 矩阵乘，每一步都有可参考的实现。

仓库链接：https://atomgit.com/cann/asc-devkit