昇腾 Ascend C 算子开发编程语言全解析：从并行计算原理到 SPMD 与流水线编程模型

本篇技术指南以昇腾算子开发编程语言 Ascend C 为核心，系统讲解其背后的并行计算基本原理、AI Core 硬件抽象、SPMD 编程模型与流水线编程范式，并延伸至语法扩展、API 体系与向量算子编程实例。读者学完后，将能够理解 Ascend C 为何这样设计、算子如何被切分并在多核上并行调度，以及如何借助 TQue/TPipe/TBuf 完成数据通路搭建，从而具备阅读与编写自定义算子的基础能力。本文基于仓库 03Compiler/07CANN 下的系列文档整理而成，是深入 CANN 算子开发体系的入门与进阶参考。

并行计算基本原理

并行计算是一种计算模式，它涉及同时执行多个计算任务或同时执行多个进程。与之相对，串行计算按顺序执行一个任务，然后再执行下一个任务。并行计算的本质，是让多个任务或进程同时执行，以提高整体计算性能和效率。

并行计算可以在多个硬件处理单元（如多个处理器、多个加速硬件、多个计算节点等）上同时执行任务。这里的加速硬件既可以是昇腾 AI 处理器，也可以是 GPU、TPU 和 FPGA 等设备。并行硬件有助于处理大规模的计算密集型问题，加快计算速度，提高系统的吞吐量，广泛应用于科学研究、工程设计、图形处理、数据分析等领域。

并行体系结构分类

从计算机硬件、系统以及应用三个层面，大致包含三类并行体系结构：指令级并行（ILP）、线程级并行（TLP） 以及请求级并行（RLP）。

指令级并行（Instruction-Level Parallelism，ILP） 指处理器内部多条机器指令能够在同一时钟周期执行。这种并行性不需要程序员手动优化代码，可由处理器自身的硬件管理。提升指令级并行的主要技术包括：

• 超标量架构（Superscalar）：允许每个时钟周期发射多条指令到不同的执行单元。超标量处理器具备多个执行单元（如整数运算、浮点运算、加载/存储等），可以同时执行多个操作；
• 流水线（Pipelining）：将指令分解为小步骤，每个步骤由不同的处理器部件顺序完成，一个指令的各个阶段可以与其他指令的阶段重叠，从而同一时刻有多个指令处于不同的执行阶段。例如，一个典型的四段流水线包括取指（IF）、译码（ID）、执行（EXE）、写回（WB）等阶段。

线程级并行（Thread-Level Parallelism，TLP） 通过创建多个线程实现并行计算。在多核和多处理器系统中，这些线程可以真正并行地运行在不同的处理器或核心上。线程级并行常见于操作系统、数据库系统以及服务端应用等领域，通常需要程序员显式地通过编程创建和管理线程。

请求级并行（Request-Level Parallelism，RLP） 通常出现在应用服务中，例如当多个独立的客户端发送请求到服务器时，服务器会创建不同的处理流程同时处理这些请求。每个请求可能涉及不同的资源和计算路径，因此可以并行处理，从而提高服务的执行能力和响应速度。

从软件设计和编程模型的角度看，并行体系结构又可划分为数据级并行和任务级并行：

• 数据级并行（Data-Level Parallelism）：将较大的数据块分割成较小的块，然后在多个处理单元上并行处理。每个处理单元运行相同的操作，但作用于不同的数据片段。数据并行特别适合数组、向量和矩阵等数据结构，常见于科学计算和图像处理等领域；
• 任务级并行（Task-Level Parallelism）：将工作分解成独立的任务，这些任务可以同时在不同的处理单元上执行。任务之间可能互相依赖，也可能完全独立，通常需要程序员设计能有效利用并行硬件特性的算法和程序结构，广泛用于软件工程、复杂事件处理和多媒体应用。

一个直观的例子：大型企业发工资。 假设一个大型企业每月需要发放数百万名员工的工资。从数据级并行的角度看，全部员工的工资计算过程被分割成小块，每块包含一部分员工的工资数据，然后分发到不同处理器上；每个处理器执行完全相同的计算任务（税务扣除、福利计算等），但仅处理分配到的数据块，例如处理器 A 负责员工列表 1～1000 的工资，处理器 B 负责员工列表 1001～2000 的工资，各自独立生成工资条。而从任务级并行的角度看，各个处理器可能负责不同的计算任务，例如处理器 A 负责计算所有员工的税务扣除，处理器 B 负责计算所有员工的福利待遇；此时每个处理器同时读取全部的工资数据集，仅对数据执行特定任务，最后合并所有处理器的输出生成最终工资条。

此外，迈克尔·费林（Michael Flynn）于 1972 年提出费林分类法，根据计算机体系结构中指令流和数据流的组织方式，将计算机系统划分为四类：

• 单指令流单数据流（SISD）：顺序执行一串指令，每个指令作用于单个数据元素。这类系统代表单处理器系统，任何时刻只有一条指令在执行，并且只操作一份数据。大多数早期计算机和现代非并行处理器均遵循 SISD 架构；
• 单指令流多数据流（SIMD）：允许不同处理器同时对多个数据元素执行同一条指令操作，适合数据并行任务（图像视频处理、矩阵运算等）。向量处理器和图形处理单元（GPU）可归入 SIMD 架构范畴；
• 多指令流单数据流（MISD）：同一时间有多条指令操作同一个数据流，实际应用较少，使用场景相对特殊。理论上的例子是容错计算，即多个处理单元对相同数据执行不同操作以检测错误；
• 多指令流多数据流（MIMD）：支持多个独立的指令流，每个都操作不同的数据流，即每个处理器可以执行不同任务、针对不同数据集。超标量处理器、大多数现代多核与多处理器系统均属于 MIMD 架构；昇腾 AI 处理器的向量计算单元与矩阵计算单元也可以视为 MIMD 架构。MIMD 非常适合需要复杂任务协调的问题，如服务器处理多个用户请求、超级计算机执行科学模拟。

费林分类法有助于简化并行计算的理解，但随着 GPU 与众核处理器的出现以及异构计算的流行，该分类已不能完全覆盖所有并行计算模式，因此引入了更复杂的并行处理模式。例如单程序多数据（Single Program Multiple Data，SPMD），它在 SIMD 基础上做了扩展，属于并行计算的编程模型：当硬件的各处理器有自己独立控制部件时，可通过软件编程让各处理器并行地执行同一个程序，但每个处理器处理不同的数据。Ascend C 正是基于昇腾 AI Core 形成了 SPMD 编程模型，这一点将在后文详细展开。

大模型并行原理

随着深度学习技术的快速进步，模型的规模和复杂性正以前所未有的速度增长。语言大模型（LLM）与视觉大模型（VLM）在提升任务性能方面取得了显著成果，但同时也带来了巨大的计算挑战——例如 OpenAI 的 ChatGPT 拥有 1750 亿个参数。传统的单机训练方法在处理这类庞大模型时，会遇到内存限制、计算速度过慢和训练时间过长等难以逾越的问题。大规模并行计算技术因此成为提高大规模模型训练效率和性能的关键。这里阐述的大模型并行加速策略，结合 Ascend C 算子开发语言，是华为昇腾大模型解决方案的基石。

目前针对 AI 大模型的并行化训练主要有两类常用方式：数据并行和模型并行。

数据并行将大规模数据集划分为多个批次（batch），分配给不同的计算节点并行处理。每个计算节点都有一个模型副本，独立执行前向和后向传播并计算梯度；随后所有计算节点通过通信协议交换梯度信息并汇总，通过聚合得到的全局梯度更新模型参数，确保所有节点上的模型副本保持同步。数据并行使模型训练可以扩展到更多的数据和计算资源，从而加速训练过程。

模型并行用于单个模型太大而无法放入单个计算节点内存的场景：将模型的不同部分（如不同的层或子网络）分布到不同节点上。各分布式节点负责模型的一部分计算，并可能频繁进行跨节点通信以同步中间状态和梯度信息。模型并行允许处理更大的模型，但由于通信开销较大，如何设计和优化模型的分割与通信策略显得尤为重要。

模型并行中存在一种特殊形式——张量并行（Tensor Parallelism），它专注于模型中单个层的并行。例如，对于一个巨大的矩阵乘法操作，张量并行会将矩阵分割成更小的块，在不同计算节点上执行这些较小规模的乘法操作，并要求跨节点协调执行和数据交换以完成整个层的运算。张量并行常用于全连接层、卷积层等参数量大的层，可减少单个节点所需处理的参数数量，从而克服内存限制问题。

并行效率量化原理

阿姆达尔定律（Amdahl's Law） 由吉恩·阿姆达尔在 1967 年提出，用于估计程序在并行化后的理论性能提升。该定律指出，一个程序的加速比上限受其串行部分比例的限制，公式如下：

其中，加速比 S 表示加速后总体性能的提升倍数，p 是程序中可以并行化的代码部分所占的比例（介于 0 和 1 之间），N 是用于并行处理的处理器数量。

阿姆达尔定律的核心观点是：并行计算的最大性能提升受限于程序中无法并行化的部分。根据公式，即使并行部分的速度被无限加速（即 N → ∞），总体加速比 S 也永远不会超过 1/(1-p)。例如，如果一个程序有 10% 的代码是串行的（p=0.9），那么即使在无限多的处理器上运行，最大加速比也只能达到 10 倍，因为串行部分的执行时间成了瓶颈。

阿姆达尔定律揭示了优化程序性能的重要策略：尽可能增加程序的可并行化比例。同时它也揭示了并行计算面临的挑战，特别是对于那些难以大幅度并行化的应用或算法。实际应用中，程序员和系统设计师会使用阿姆达尔定律评估并行化的潜在价值，并决定在硬件和软件层面需要投入多少资源进行并行优化。

Ascend C 编程体系

算子是编程和数学中的重要概念，是用于执行特定操作的符号或函数，处理输入值并生成输出值。一元算子对单个操作数进行操作（如取反或递增），二元算子对两个操作数执行操作（如加法或赋值），关系算子用于比较值之间的关系，逻辑算子用于在逻辑表达式中组合条件，位运算符操作二进制位，赋值算子将值分配给变量。算子在编程语言中定义了基本操作，在数学中描述了对数学对象的操作（如微积分中的导数和积分算子）。

面向算子开发场景的编程语言 Ascend C，原生支持 C 和 C++ 标准规范，最大化匹配用户的开发习惯；通过多层接口抽象、自动并行计算、孪生调试等关键技术，极大提高算子开发效率，助力 AI 开发者低成本完成算子开发和模型调优部署。在 CANN 异构计算架构 中，Ascend C 正是 CANN 面向算子开发场景开放的低阶 API 编程语言，与 AscendCL、GE 图引擎、HCCL 集合通信库等组件协同构成完整的算子开发与执行生态。

使用 Ascend C 进行自定义算子开发的突出优势有：

• C/C++ 原语编程：原生支持 C 和 C++ 标准规范，最大化匹配用户开发习惯；
• 编程模型屏蔽硬件差异：编程范式提高开发效率；
• 类库 API 封装：从简单到灵活，兼顾易用与高效；
• 孪生调试：CPU 侧模拟 NPU 侧的行为，可优先在 CPU 侧调试。

Ascend C 的编程模型

AI Core 是昇腾 AI 处理器执行计算密集型任务的真正算力担当。理解 AI Core 的硬件抽象（计算、存储和 DMA 搬运单元），以及 SPMD 模型和流水线编程范式，是有效提高 Ascend C 算子并行执行效率、实现更高加速比计算任务的关键。

AI Core 硬件抽象

使用 Ascend C 编程语言开发的算子运行在 AI Core 上。AI Core 是昇腾 AI 处理器中的计算核心，一个 AI 处理器内部有多个 AI Core。AI Core 包含计算单元、存储单元、搬运单元、控制单元等核心部件。为了屏蔽不同 AI Core 硬件资源上可能存在的差异性，Ascend C 对其进行了统一的硬件抽象表示：

• 计算单元包括三种基础计算资源：矩阵计算单元（Cube）、向量计算单元（Vector） 和标量计算单元（Scalar）；
• 存储单元即 AI Core 的内部存储，统称为本地内存（Local Memory）；与之对应，AI Core 的外部存储称为全局内存（Global Memory）；
• DMA 搬运单元负责在全局内存和本地内存之间搬运数据。

AI Core 内部的异步并行计算过程如下：标量计算单元读取指令序列，并把向量计算、矩阵计算、数据搬运指令发射给对应单元的指令队列；向量计算单元、矩阵计算单元、数据搬运单元异步地并行执行接收到的指令（上图中蓝色箭头所示的指令流）。由于不同指令间可能存在依赖关系，为了保证不同指令队列间的指令按照正确的逻辑关系执行，标量计算单元也会给对应单元下发同步指令（上图中绿色箭头所示的同步信号流）。

AI Core 内部数据处理的基本过程为：DMA 搬入单元把数据搬运到 Local Memory → Vector/Cube 计算单元完成数据计算并把结果写回 Local Memory → DMA 搬出单元把处理好的数据搬运回 Global Memory（上图中红色箭头所示的数据流）。

SPMD 编程模型

SPMD（Single-Program Multiple-Data）数据并行是典型的并行计算方法。其原理如下：假设从输入数据到输出数据需要经过 3 个阶段任务的处理（T1、T2、T3），SPMD 会启动一组进程并行处理待处理的数据：对待处理数据切分，把切分后的数据分片分发给不同进程，每个进程对自己的数据分片执行 3 个任务的处理。

Ascend C 算子编程即 SPMD 编程：将需要处理的数据拆分，并同时在多个计算核心上运行以获取更高性能。多个 AI Core 共享相同的指令代码，每个核上的运行实例唯一的区别是 block_idx 不同，每个核通过不同的 block_idx 识别自己的身份。block 的概念类似于 SPMD 中进程的概念，它与物理概念上芯片固定的核数不同，可以自行定义设置；block_idx 就是标识进程唯一性的进程 ID。一般而言，为了充分利用计算资源，block 数量设置为物理核数的整数倍。

流水线编程范式

在计算机编程方法中，做到软件与硬件的结合对性能提升有很大帮助，其中很关键的一点是：在代码流执行过程中让所有计算资源（硬件）都处于高占用率状态并进行有效的运算，避免长时间空闲。将程序流水化是达到上述效果的最好方法——把完整的任务进行模块化处理，多个模块之间形成流水关系，并使用队列来处理不同模块之间的异步并行。

Ascend C 编程范式是一种流水线式的编程范式：把算子核内的处理程序分成多个流水任务，通过队列（TQue）完成任务间通信和同步，并通过统一的内存管理模块（TPipe）管理任务间通信内存。流水编程范式的关键是流水任务设计——流水任务指单核处理程序中主程序调度的并行任务，在核函数内部可通过流水任务实现数据的并行处理，进一步提升性能。

以下图为例，单核处理程序的功能被拆分成 3 个流水任务：阶段 1、阶段 2、阶段 3，每个任务专注于完成单一功能；待处理数据被切分成 n 片（数据分块 1~n），每个任务需要依次完成 n 个数据切片的处理。阶段间的箭头表达数据间的依赖关系，比如阶段 1 处理完数据分块 1 之后，阶段 2 才能对数据分块 1 进行处理。

若 n=3，即待处理数据被切分成 3 片，流水任务运行起来如下图所示。从运行图中可以看出：对于同一片数据，阶段 1、阶段 2、阶段 3 之间的处理具有依赖关系，需要串行处理；不同的数据切片，在同一时间点可以有多个任务并行处理，由此达到任务并行、提升性能的目的。

Ascend C 提供了一种通信队列实现各阶段间的数据通信与同步，即 TQue。TQue 的分类大致有两种：VECIN 和 VECOUT，对应计算前后数据的搬入和搬出。以向量算子为例：阶段 1 为数据搬入阶段，阶段 2 为向量计算阶段，阶段 3 为数据搬出阶段，三者调用不同的芯片内单元，从而实现异步并行。对于特定的一块数据分块：在阶段 1 的最后，数据被放入 VECIN 队列中；随后在阶段 2 被取出，执行相关操作后放入 VECOUT 队列；最后在阶段 3 被从 VECOUT 队列中取出并搬出。由于不同算子需求不同，每个数据分块需要经过的阶段数量和内容也有所差异。

Ascend C 语法扩展与编程对象

Ascend C 的本质构成其实是标准 C++ 加上一组扩展的语法和 API（详见 Ascend C 语法扩展）。它采用华为自研的毕昇编译器，设备侧编程采用 C/C++ 语法扩展，允许函数执行空间和地址空间作为合法的类型限定符，提供在主机（Host）侧和设备（Device）侧独立执行的能力，以及针对不同地址空间的访问能力。

函数执行限定符

函数执行空间限定符指示函数是在主机上执行还是在设备上执行，以及它是否可从主机或设备调用，主要有三种：

• __global__ 执行空间限定符：声明一个 kernel 函数（Ascend C 算子的入口函数），具有如下性质：在设备上执行；只能被主机侧函数调用；__global__ 只是表示这是设备侧函数的入口，并不表示具体的设备类型；一个 __global__ 函数必须返回 void 类型，并且不能是 class 的成员函数；主机侧调用 __global__ 函数必须使用 <<<>>>（内核调用符）异构调用语法；__global__ 的调用是异步的，函数返回并不表示 kernel 函数在设备侧已经执行完成，如需同步须使用运行时提供的同步接口（如 aclrtSynchronizeStream(...)）显式同步；
• __aicore__ 执行空间限定符：声明的函数在 AI Core 设备上执行；只能被 __global__ 函数或其他 AI Core 函数调用；
• __host__ 执行空间限定符：声明的函数（通常不显式声明）只能在主机侧执行，只能被主机侧函数调用；__global__ 和 __host__ 不能一起使用。

典型使用示例：

// 定义 aicore 函数
__aicore__ void bar() {}

// 定义核函数
__global__ __aicore__ void foo() { bar(); }

// 定义 Host 函数
int main() {}

地址空间限定符

地址空间限定符可以在变量声明中使用，指定对象分配的区域。AI Core 具备多级独立片上存储，各个地址空间独立编址，具备各自的访存指令。如果对象的类型被地址空间名称限定，则该对象将被分配在指定的地址空间中；对于指针，指向的类型也可以通过地址空间限定，指示所指向对象所在的地址空间。

• private 地址空间：大多数变量的默认地址空间，特别是局部变量，代码中不显式标识所在局部地址空间类型；
• __gm__ 地址空间：__gm__ 地址空间限定符用来表示分配于设备侧全局内存的对象，全局内存对象可以声明为标量、用户自定义结构体的指针。

Ascend C API 体系

Ascend C 算子采用标准 C++ 语法和一组编程类库 API 进行编程。Ascend C API 的操作数都是 Tensor 类型——GlobalTensor（外部数据存储空间）和 LocalTensor（核上内存空间），类库 API 分为高阶 API 和基础 API（API 全景图见 04Grmmar.md 中的编程类库 API 示意图）。

基础 API 实现基础功能，包括计算类、数据搬运、内存管理和任务同步等。使用基础 API 自由度更高，可以通过 API 组合实现自己的算子逻辑，是对计算能力的表达：

• 计算类 API：包括标量计算 API、向量计算 API、矩阵计算 API，分别调用标量计算单元、向量计算单元、矩阵计算单元执行计算；
• 数据搬运 API：计算 API 基于本地内存数据进行计算，数据需要先从全局内存搬运至本地内存，计算完成后再搬出至全局内存。执行搬运过程的接口称为数据搬运接口，如 DataCopy 接口；
• 内存管理 API：用于分配板上管理内存，如 AllocTensor、FreeTensor 接口。由于板上内存较小，通常无法存储完整数据，因此采用动态内存管理实现板上内存复用；
• 任务同步 API：完成任务间的通信和同步，如 EnQue、DeQue 接口。不同指令异步并行执行，为保证不同指令队列间的指令按正确逻辑关系执行，需要向不同组件发送同步指令；任务同步类 API 内部即完成该过程，开发者无需关注内部实现。

基础 API 中的计算 API 按数据操作方法不同分为几种计算方式：

命名	说明
整个 Tensor 参与计算	通过运算符重载的方式实现，支持 +、-、*、/、|、&、<、>、<=、>=、==、!=，实现计算的简化表达，例如 dst = src1 + src2
Tensor 前 n 个数据计算	针对源操作数的连续 n 个数据进行计算并连续写入目的操作数，解决一维 Tensor 的连续计算问题，例如 Add(dst, src1, src2, n)；这是一般开发过程中使用最频繁的接口，兼具功能性与易用性
Tensor 高维切分计算	功能灵活的计算 API，充分发挥硬件优势，支持对每个操作数的 Repeat times（迭代次数）、Block stride（单次迭代内不同 block 间地址步长）、Repeat stride（相邻迭代间相同 block 的地址步长）、Mask（控制参与运算的计算单元）进行操作，操作单元最小，可针对不同步长实现最细致的操作

高阶 API 封装常用算法逻辑，如 Matmul、Softmax 等，可减少重复开发、提高开发效率，是对某种特定算法的表达。例如使用高阶 API 完成 Matmul 算子时，需要创建一个矩阵乘法类进行运算，入参包含两个相乘矩阵（A 矩阵与 B 矩阵）信息、输出结果矩阵（C 矩阵）信息、矩阵乘偏置（Bias）信息，上述信息中包括对应矩阵数据的内存逻辑位置、数据存储格式、数据类型、转置使能参数。创建矩阵乘法类后，可直接完成对矩阵 A、B 和 Bias 的设置、矩阵乘法操作及结果输出，无需自主实现复杂的数据通路和运算操作。

数据存储：GlobalTensor 与 LocalTensor

根据 Ascend C 对 AI Core 的硬件抽象设计，AI Core 内部的存储统一用 Local Memory 表示，外部的存储统一用 Global Memory 表示。Ascend C 使用 GlobalTensor 作为 Global Memory 的数据基本操作单元，用 LocalTensor 作为 Local Memory 的数据基本操作单元。Tensor（张量）是各种指令 API 直接处理的对象，也是数据的载体。

外部存储数据空间（GlobalTensor） 存放 AI Core 外部存储的全局数据，其原型定义如下：

template <typename T> class GlobalTensor {
    void SetGlobalBuffer(__gm__ T* buffer, uint32_t bufferSize);
    const __gm__ T* GetPhyAddr();
    uint64_t GetSize();
    GlobalTensor operator[](const uint64_t offset);
}

• SetGlobalBuffer：传入全局数据的指针，并手动设置 buffer size，初始化 GlobalTensor；
• GetPhyAddr：返回全局数据的地址，类型 T 支持所有数据类型，但需遵循使用此 GlobalTensor 的指令的数据类型支持情况；
• GetSize：返回 Tensor 中的 element 个数；
• operator[]：指定偏移返回一个 GlobalTensor，offset 单位为 element。

核内数据存储空间（LocalTensor） 存放 AI Core 内部存储的数据，其原型定义如下：

template <typename T> class LocalTensor {
    T GetValue(const uint32_t offset) const;
    template <typename T1> void SetValue(const uint32_t offset, const T1 value) const;
    LocalTensor operator[](const uint32_t offset) const;
    uint32_t GetSize() const;
    void SetUserTag(const TTagType tag);
    TTagType GetUserTag() const;
}

• GetValue：获取 LocalTensor 中的某个值，返回 T 类型的立即数；
• SetValue：设置 LocalTensor 中的某个值，offset 单位为 element；
• operator[]：获取距原 LocalTensor 起始地址偏移量为 offset 的新 LocalTensor，offset 不能超过原有 LocalTensor 的 size 大小，单位为 element；
• GetSize：获取当前 LocalTensor 尺寸大小；
• SetUserTag / GetUserTag：开发者自定义设置/获取 Tag 信息。

任务通信与同步：TQue 与 TPosition

Ascend C 使用 TQue 队列完成任务之间的数据通信和同步。TQue 队列管理不同层级的物理内存时，用一种抽象的逻辑位置（TPosition）表达各级别的存储，代替了片上物理存储的概念，开发者无需感知硬件架构。

TPosition 类型包括：VECIN、VECCALC、VECOUT、A1、A2、B1、B2、CO1、CO2。其中 VECIN、VECCALC、VECOUT 主要用于向量编程；A1、A2、B1、B2、CO1、CO2 用于矩阵编程。各枚举值含义如下：

TPosition	具体含义
GM	全局内存（Global Memory），对应 AI Core 的外部存储
VECIN	用于向量计算，搬入数据的存放位置，在数据搬入矩阵计算单元时使用此位置
VECOUT	用于向量计算，搬出数据的存放位置，在将矩阵计算单元结果搬出时使用此位置
VECCALC	用于向量计算/矩阵计算，在计算需要临时变量时使用此位置
A1	用于矩阵计算，存放整块 A 矩阵，可类比 CPU 多级缓存中的二级缓存
B1	用于矩阵计算，存放整块 B 矩阵，可类比 CPU 多级缓存中的二级缓存
A2	用于矩阵计算，存放切分后的小块 A 矩阵，可类比 CPU 多级缓存中的一级缓存
B2	用于矩阵计算，存放切分后的小块 B 矩阵，可类比 CPU 多级缓存中的一级缓存
CO1	用于矩阵计算，存放小块结果 C 矩阵，可理解为 Cube Out
CO2	用于矩阵计算，存放整块结果 C 矩阵，可理解为 Cube Out

一个使用 TQue 数据结构的样例程序：

TQue<TPosition::VECIN, BUFFER_NUM> que;
LocalTensor<half> tensor1 = que.AllocTensor();
que.FreeTensor<half>(tensor1);
que.EnQue(tensor1);
LocalTensor<half> tensor1 = que.DeQue<half>();

• 第 1 行：创建队列，VECIN 是 TPosition 枚举中的一员，BUFFER_NUM 参数是队列深度，que 是自定义队列名称；
• 第 2 行：使用 AllocTensor() 在片上分配空间给一个 LocalTensor，分配的默认大小为 que 初始化时设置的一块的大小，也可手动设置大小但不能超过分配大小；
• 第 3 行：使用 FreeTensor() 释放 Tensor，需与 AllocTensor() 成对使用；
• 第 4 行：将 LocalTensor 加入指定队列，利用队列进行不同任务间的数据同步与通信；
• 第 5 行：将 LocalTensor 从队列中取出，以便进行后续计算等操作。

资源管理：TPipe 与 double buffer

在 Ascend C 中，流水任务间数据传递使用到的内存统一由资源管理模块 TPipe 管理。TPipe 作为片上内存管理者，通过 InitBuffer 接口对外提供 TQue 内存初始化功能，开发者可通过该接口为指定的 TQue 分配内存。TQue 队列内存初始化完成后，需要内存时调用 AllocTensor 为 LocalTensor 分配内存；当 LocalTensor 完成计算不再使用时，再调用 FreeTensor 回收内存。

一个使用 TPipe 的样例：

TPipe pipe;
pipe.InitBuffer(que, num, len);

其中 que 为指定的已创建队列名称；num 为分配的内存块数，num=2 时开启 double buffer 优化；len 为分配的一块内存大小（单位 Bytes）。

double buffer 优化机制：执行于 AI Core 上的指令队列主要包括 Vector 指令队列（V）、Matrix 指令队列（M）和存储移动指令队列（MTE2、MTE3）。不同指令队列间的相互独立性和可并行执行特性是 double buffer 优化机制的基石。double buffer 基于 MTE 指令队列与 Vector 指令队列的独立性和可并行性，通过将数据搬运与 Vector 计算并行执行，隐藏数据搬运时间并降低 Vector 指令的等待时间，最终提高 Vector 单元的利用效率。

临时变量：TBuf

使用 Ascend C 编程时可能会用到一些临时变量，例如在 Compute 阶段使用复杂的数据结构。这些临时变量占用的内存可以使用 TBuf 管理，存储位置通过模板参数设置，可以设置为不同的 TPosition 逻辑位置。建议将临时变量初始化成为算子类成员之一，避免重复申请与释放，从而提升算子性能。

TBuf 占用的存储空间通过 TPipe 管理，可通过 InitBuffer 接口为其进行内存初始化，之后通过 Get 获取指定长度的 Tensor 参与计算。与 TQue 分配内存相比，TBuf 有以下差异：

• 为 TBuf 分配的内存空间只能参与计算，无法执行 TQue 队列的入队出队操作；
• 调用一次内存初始化接口，TPipe 只为 TBuf 分配一块内存；TQue 队列可以通过参数设置申请多块内存。若要使用多个临时变量，需要定义多个 TBuf 数据结构并分别调用 InitBuffer 进行内存初始化；
• 使用 TBuf 时可以不重复进行申请释放内存操作。

一个使用 TBuf 的样例：

TBuf<TPosition::pos> calcBuf;
pipe.InitBuffer(calcBuf, len);
LocalTensor<half> temtensor1 = calcBuf.Get<half>();
LocalTensor<half> temtensor2 = calcBuf.Get<half>(128);

• 第 1 行：声明临时变量，pos 为队列逻辑位置 TPosition，可以为 VECIN、VECCALC、VECOUT、A1、A2、B1、B2、CO1、CO2；
• 第 2 行：使用 TPipe 的 InitBuffer 操作对临时变量进行初始化，只需声明名称和长度 len（单位 Bytes）；
• 第 3、4 行：分配临时的 LocalTensor，使用 TBuf 的 Get() 方法；不传入参时分配空间大小为初始化时 len 的大小，传入参时可以指定长度分配临时 Tensor 大小，但长度不能超过 len。

向量编程范式：以 add_custom 算子为例

基于 Ascend C 编程范式实现自定义向量算子的流程（详见 Ascend C 编程范式）由三个步骤组成：算子分析是编程的前置任务，负责明确自定义算子的各项需求（输入输出、使用的 API 接口等）；核函数的定义和封装是编程的第一步，负责声明核函数名称并提供进入核函数运算逻辑的接口；基于算子需求实现算子类是整个核函数的核心计算逻辑，由内存初始化、数据搬入、算子计算逻辑实现、数据搬出四部分组成，后三者又被称为算子的实现流程。

自定义向量算子核心部分一般由两个函数组成：Init() 函数（初始化函数）与 Process() 函数（执行函数）。Init() 完成板外数据定位以及板上内存初始化；Process() 完成向量算子的实现，分成三个流水任务：CopyIn、Compute、CopyOut——CopyIn 负责板外数据搬入，Compute 负责向量计算，CopyOut 负责板上数据搬出。流水线任务之间存在数据依赖，需要通过 TQue 队列（提供 EnQue、DeQue 等基础 API）完成数据通信与同步。在向量编程模型中，使用到的 TQue 类型为：搬入数据的存放位置 VECIN、搬出数据的存放位置 VECOUT。

下面以 add_custom 这一基本的向量算子为例，逐步介绍向量编程范式。

算子分析

• 明确算子的数学表达式：Ascend C 提供的向量计算接口的操作元素都为 LocalTensor，输入数据需先搬运进片上存储。以 Add 算子为例，数学表达式为 z = x + y，使用计算接口完成两个输入参数相加，得到最终结果后再搬出到外部存储；
• 明确输入和输出：Add 算子有两个输入 x 与 y，输出为 z。本样例中算子输入支持的数据类型为 half(float16)，输出数据类型与输入相同；输入支持 shape（8，2048），输出 shape 与输入相同；支持的数据格式（shape）为 ND；
• 确定算子实现所需接口：使用 DataCopy 实现数据搬移；向量计算实现较为简单，使用基础 API 完成计算逻辑，加法算子中使用双目指令接口 Add 实现 x+y；使用 EnQue、DeQue 等接口对 TQue 队列进行管理。

核函数定义与封装

本样例中函数原型名为 add_custom，有 3 个参数 x、y、z，其中 x、y 为输入内存，z 为输出内存。根据核函数定义规则，使用 __global__ 函数类型限定符标识它是一个核函数，可以被 <<<...>>> 调用；使用 __aicore__ 函数类型限定符标识该核函数在设备端 AI Core 上执行；为方便起见，统一使用 GM_ADDR 宏修饰入参，表示其为入参在内存中的位置。函数原型中首先实例化对应的算子类，并调用该算子类的 Init() 和 Process() 函数（Init() 负责内存初始化，Process() 负责算子实现核心逻辑）。对核函数的调用进行封装得到 add_custom_do 函数便于主程序调用——该封装函数仅在编译运行 NPU 侧算子时用到，编译运行 CPU 侧算子时可直接调用 add_custom 函数。调用核函数时，除传入参数 x、y、z 外，还需要使用 <<<…>>> 传入 blockDim（核函数执行的核数）、l2ctrl（保留参数，设置为 nullptr）、stream（应用程序中维护异步操作执行顺序的任务流对象）来规定核函数的执行配置：

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z){
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

#ifndef __CCE_KT_TEST__

// call of kernel function
void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z){
    add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}

#endif

算子数据通路

向量算子三阶段任务流水的数据通路可纵向分为两部分：上部分为发生在外部存储（Global Memory）中的数据流通过程，下部分为发生在 AI Core 内（Local Memory）中的数据流通过程；横向分为三部分，指代 CopyIn、Compute 和 CopyOut 三个阶段中的数据流通过程。发生在 AI Core 内的任务间数据传递统一由 TPipe 资源管理模块管理。

• CopyIn 任务：先将执行计算的数据 xGm、yGm 从外部存储通过 DataCopy 接口传入板上，存储为 xLocal、yLocal，并通过 EnQue 接口传入数据搬入队列 inQueueX、inQueueY 中，以便进行流水模块间的数据通信与同步；
• Compute 任务：先将 xLocal、yLocal 使用 DeQue 接口从数据搬入队列中取出，使用相应向量运算 API 执行计算得到结果 zLocal，再将 zLocal 通过 EnQue 接口传入数据搬出队列 outQueueZ 中；
• CopyOut 任务：先将结果数据 zLocal 使用 DeQue 接口从数据搬出队列中取出，使用 DataCopy 接口将板上数据传出到外部存储 zGm 中。

算子类实现

算子类成员定义：在算子类中，需要声明对外开放的内存初始化函数 Init() 和核心处理函数 Process()。为实现向量算子核内计算流水操作，将 Process() 分为三个部分：数据搬入阶段 CopyIn()、计算阶段 Compute() 与数据搬出阶段 CopyOut() 三个私有类成员。此外还依次声明内存管理对象 pipe、输入数据 TQue 队列管理对象 inQueueX 和 inQueueY、输出数据 TQue 队列管理对象 outQueueZ 以及管理输入输出 Global Memory 内存地址的对象 xGm、yGm 与 zGm，均作为私有成员使用：

class KernelAdd {

public:
    __aicore__ inline KernelAdd() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){}
    __aicore__ inline void Process(){}

private:
    __aicore__ inline void CopyIn(int32_t progress){}
    __aicore__ inline void Compute(int32_t progress){}
    __aicore__ inline void CopyOut(int32_t progress){}

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

初始化函数 Init() 实现：在多核并行计算中，每个核计算的数据是全部数据的一部分，Ascend C 核函数是单个核的处理函数，因此需要获取每个核负责的对应位置的数据，并为声明的输入输出 TQue 队列分配相应内存空间。Init() 通过计算得到该核所负责的数据所在位置：x、y、z 表示 3 个入参在片外的起始地址；BLOCK_LENGTH 表示单个核负责的数据长度，为数据全长与参与计算核数的商；GetBlockIdx() 是与硬件感知相关的 API 接口，可得到核所对应的编号（该样例中为 0-7）。通过这种方式得到该核函数需要处理的输入输出在 Global Memory 上的内存偏移地址，并设置在 Global Tensor 中；随后通过 TPipe 内存管理对象调用 InitBuffer() 为输入输出 TQue 分配内存，接口入参依次为 TQue 队列名、是否启动 double buffer 机制以及单个数据块的大小（而非长度）：

__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    xGm.SetGlobalBuffer((__gm__ half*)x + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__ half*)y + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__ half*)z + BLOCK_LENGTH * GetBlockIdx(), BLOCK_LENGTH);
    pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));
    pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half));
}

核心处理函数 Process() 实现：基于向量编程范式，将核函数的实现分为 3 个基本任务 CopyIn、Compute、CopyOut。考虑到每个核内的数据仍被进一步切分成小块，需要循环执行上述步骤以得到最终结果。Process() 首先计算每个核内的分块数量，确定循环执行三段流水任务的次数，随后依此循环顺序执行 CopyIn()、Compute() 和 CopyOut()：

__aicore__ inline void Process()
{
    constexpr int32_t loopCount = TILE_NUM * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++) {
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

__aicore__ inline void CopyIn(int32_t progress)
{
    LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
    LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();

    DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
    DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);

    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}

__aicore__ inline void Compute(int32_t progress)
{
    LocalTensor<half> xLocal = inQueueX.DeQue<half>();
    LocalTensor<half> yLocal = inQueueY.DeQue<half>();
    LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();

    Add(zLocal, xLocal, yLocal, TILE_LENGTH);
    outQueueZ.EnQue<half>(zLocal);

    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

__aicore__ inline void CopyOut(int32_t progress)
{
    LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
    DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
    outQueueZ.FreeTensor(zLocal);
}

• CopyIn() 实现：使用 AllocTensor 接口为参与计算的输入分配板上存储空间（入参数据类型为 half，分配空间大小也为 half）；使用 DataCopy 接口将 GlobalTensor 数据拷贝到 LocalTensor（xGm、yGm 存储该核所需处理的所有输入，根据分块对应编号找到相关分块数据拷贝至板上）；使用 EnQue 将 LocalTensor 放入 VECIN 的 TQue 中；
• Compute() 实现：使用 DeQue 从 VECIN 中取出输入 x 和 y；使用 AllocTensor 为输出分配板上存储空间；使用 Ascend C 接口 Add 完成向量计算——该接口是双目指令 2 级接口，入参分别为目的操作数、源操作数 1、源操作数 2 和输入元素个数；使用 EnQue 将计算结果 LocalTensor 放入 VECOUT 的 TQue 中；使用 FreeTensor 释放不再使用的两个输入 LocalTensor；
• CopyOut() 实现：使用 DeQue 从 VECOUT 的 TQue 中取出目标结果 z；使用 DataCopy 将 LocalTensor 数据拷贝到 GlobalTensor 上；使用 FreeTensor 回收不再使用的 LocalTensor。

算子切分策略

Ascend C 算子编程是 SPMD 编程，使用多个核进行并行计算，在单个核内还将数据根据需求切分成若干份以降低每次计算负荷。需要注意的是，Ascend C 中涉及的核数并不是实际执行硬件中所拥有的处理器核数，而是"逻辑核"的数量，即同时运行了多少个算子实例（同时执行此算子的进程数量）。一般建议逻辑核数量是实际处理器核数的整数倍；如果条件允许，还可以进一步将每个待处理数据一分为二，开启 double buffer 机制，实现流水线间并行，进一步减少计算单元的闲置问题。

在本 add_custom 算子样例中，数据切分参数定义如下：

constexpr int32_t TOTAL_LENGTH = 8 * 2048;   // 数据全长，约束输入数据长度
constexpr int32_t USE_CORE_NUM = 8;          // 参与计算任务的核数
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // 单个核负责计算的数据长度
constexpr int32_t TILE_NUM = 8;              // 单个核中数据的切分块数
constexpr int32_t BUFFER_NUM = 2;            // 是否开启 double buffer，不开启则 BUFFER_NUM = 1
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // 单个数据块的数据长度

具体切分过程为：设置数据整体长度 TOTAL_LENGTH 为 8×2048，平均分配到 8 个核上运行，单核处理的数据大小 BLOCK_LENGTH 为 2048；对于单核上的处理数据再进行数据切块，切分成 8 块（8 块并不意味着性能最优）；切分后的每个数据块再次切分成 2 块即可开启 double buffer，此时每个数据块的长度 TILE_LENGTH 为 128 个数据。整个切分链路为：确定一个数据的起始内存位置后，将数据整体平均分配到各个核，随后针对单核数据再次切分（8 块），并启动 double buffer 机制将每个数据块一分为二，得到单个数据块的长度 TILE_LENGTH。

小结与思考

• 并行计算基本原理：并行计算通过在多个硬件处理单元上同时执行任务来提升计算性能和效率，涵盖指令级并行（ILP）、线程级并行（TLP）和请求级并行（RLP），以及数据级并行和任务级并行，适用于科学研究、工程设计等多种领域；费林分类法（SISD/SIMD/MISD/MIMD）提供了理解并行体系结构的经典框架，SPMD 则在 SIMD 基础上扩展出更适合异构众核的编程模型；
• 并行体系结构与大模型加速：并行体系结构支持数据并行和模型并行（含张量并行）以应对 AI 大模型的计算挑战，通过切分数据或模型提升训练效率和处理能力；阿姆达尔定律则从理论上量化了并行化的收益上限，提醒开发者尽可能提升程序的可并行化比例；
• Ascend C 编程语言：Ascend C 是为昇腾 AI 处理器算子开发设计的编程语言，支持 C/C++ 标准，提供 SPMD 编程模型和流水线编程范式（CopyIn/Compute/CopyOut），通过 TQue/TPipe/TBuf 完成任务通信、资源管理与临时变量管理，通过多层接口抽象（基础 API 与高阶 API）和自动并行计算提高开发效率，助力 AI 开发者高效完成算子开发和模型调优；
• 编程范式延伸：除了向量计算范式外，还有矩阵计算范式，以及向量-矩阵混合（通常是融合算子）编程范式，感兴趣的读者可继续阅读仓库中 Ascend C 编程范式、Ascend C 语法扩展 与 昇腾异构计算架构 CANN 等系列文档。