1-一个Add算子的前世今生

1-1 AI Core 架构抽象回顾

AI Core 的逻辑组成

• AI Core 是昇腾 AI 处理器中的主要计算核心。进行 Ascend C 算子开发时，可以把 AI Core 抽象为三类核心资源：
• 计算单元：负责实际运算。
• Scalar：标量计算与控制相关处理。
• Vector：向量计算，适合逐元素、向量类操作。
• Cube：矩阵/块矩阵类计算。
• 存储单元：包括片上 Local Memory，用于保存当前 AI Core 正在处理的数据。
• 搬运单元：在 Global Memory 与 Local Memory，以及不同逻辑存储位置之间搬运数据。

• AI Core 的向量计算采用 SIMD 思想：一条指令可以并行处理一组数据，因此算子实现时需要特别关注数据块大小、对齐方式、并行切分和流水调度。

外部存储与内部存储

• Ascend C 中常用两类 Tensor 对象来表达数据所在的位置：

对象	所在位置	作用
GlobalTensor	Global Memory	表示外部存储中的全局数据，通常与核函数传入的 GM 地址关联
LocalTensor	Local Memory	表示 AI Core 片上内存中的局部数据，供计算 API 直接使用

• 典型处理流程是：

• 因此，很多 Ascend C 算子的核心都可以归纳成三个阶段：CopyIn → Compute → CopyOut。

GlobalTensor

• GlobalTensor 用于描述 Global Memory 上的数据。通常先用核函数参数获得全局内存地址，再把该地址绑定到 GlobalTensor。

• 概念上可以理解为：

• 之后可以通过下标或切片方式确定当前核需要处理的全局数据范围。

LocalTensor

• LocalTensor 用于描述 Local Memory 上的数据。Vector 等计算 API 的输入、输出通常是 LocalTensor。

• LocalTensor 可以进行偏移访问，例如从一个较大的局部 Tensor 中取得某一段数据。算子内部通常不会手工管理裸地址，而是通过 TPipe、TQue 和 LocalTensor 完成片上内存的申请、复用和释放。

逻辑位置 QuePosition

• Ascend C 使用逻辑存储位置表达不同流水阶段中的数据位置。不同逻辑位置对应不同的硬件访问通路，开发者主要通过 Queue 管理它们之间的数据流转。

• 常见位置包括：

QuePosition	典型用途
VECIN	Vector 计算输入
VECCALC	Vector 中间计算数据
VECOUT	Vector 计算输出
A1、A2、B1、B2	Cube / 矩阵计算相关输入与中间位置
CO1、CO2	Cube / 矩阵计算相关输出位置

• 使用这些逻辑位置时，不需要直接理解底层物理存储结构，但需要明确：数据处于哪个流水阶段、后续由哪个计算单元消费。

1-2 Ascend C 的编程对象

AI Core 内部并行计算的抽象

• Ascend C 算子最终运行在 AI Core 上。一个 AI 处理器中存在多个 AI Core，多核之间通过 SPMD 模型并行执行同一份核函数代码，只是不同核处理的数据区间不同。

• 单个 AI Core 内又包含计算、存储和搬运资源，因此 Ascend C 编程同时需要处理两层并行：
• 核间并行：多个 AI Core 分担不同数据。
• 核内并行：数据搬入、计算、搬出尽可能通过流水重叠执行。

SPMD 模型

• Ascend C 采用 SPMD（Single Program Multiple Data）编程模型：
• 多个 AI Core 执行同一份核函数代码。
• 每个核通过不同的 block_idx 区分自己的任务。
• 核函数中可通过 GetBlockIdx() 获取当前核的逻辑 ID。
• Host/Tiling 决定使用多少个核以及每个核处理多少数据。

• 可以把不同 AI Core 理解为执行同一个程序的多个“工作实例”：

Pipe 与 Queue

• 在核内，Ascend C 把数据处理流程拆成多个流水任务（Stage），使用：
• TPipe：管理任务间使用的片上内存。
• TQue：负责不同 Stage 之间的数据传递和同步。
• AllocTensor() / FreeTensor()：从 Queue 对应的缓冲区申请、释放 LocalTensor。
• EnQue() / DeQue()：将 Tensor 放入/取出 Queue。

• 这套机制的重点不是“保存一个容器”，而是建立 数据依赖 + 内存复用 + 流水同步。

1-3 Vector 算子开发流程——以 Add 为例

算子分析

• Add 是典型的逐元素 Vector 算子：

• 示例约定：

项目	内容
输入	x、y
输出	z
Shape	(8, 2048)
数据类型	half
Format	ND
核函数名称	add_custom

• 实现时需要明确三类接口：
• 数据搬入：使用 DataCopy 等接口将输入从 Global Memory 搬入 Local Memory。
• 计算：使用 Add 等 Vector API 完成逐元素加法。
• 数据搬出：将计算结果从 Local Memory 搬回 Global Memory。

开发流程

• Vector 算子开发可概括为：

核函数定义

• Ascend C 核函数使用 __global__ 和 __aicore__ 等限定符，设备侧代码直接在 AI Core 上执行。

• 典型形式：

• 核函数入口本身通常保持简洁：
• 创建算子实现类。
• 调用 Init() 完成地址绑定、切分参数计算和 Queue 初始化。
• 调用 Process() 执行流水。

流水任务设计

• Vector 算子通常把一次完整处理拆为三个 Stage：
1. CopyIn：Global Memory → Local Memory。
2. Compute：Local Memory 中完成 Vector 计算。
3. CopyOut：Local Memory → Global Memory。

• 当一份输入被切成多个 Progress 后，不同 Progress 可以处于不同 Stage，从而形成流水并行。例如：

• 这使搬运单元和计算单元不必长期互相等待。

算子类的典型结构

• 这套结构的意义是把内存初始化、任务调度和三个流水阶段分离，后续修改切分策略或计算逻辑时更容易维护。

Init：完成多核切分与 Queue 初始化

• 对于固定 Shape，可以在代码中预先计算：
• 总元素数 TOTAL_LENGTH。
• 使用核数 BLOCK_DIM。
• 单核处理元素数 BLOCK_LENGTH。
• 单核内部 Tile 数 TILE_NUM。
• 双缓冲数量 BUFFER_NUM = 2。
• 单个 Tile 的元素数 TILE_LENGTH。

• block_idx 决定当前核在 Global Memory 中的起始偏移：

• 之后用 SetGlobalBuffer() 绑定当前核对应的 x/y/z 数据区间，并用 pipe.InitBuffer() 为 Queue 分配 Local Memory。

Process：循环执行流水任务

• Process() 通常按 Progress 循环：

• 三个函数通过 Queue 建立数据依赖。

CopyIn

• 典型过程：
• AllocTensor() 获得 LocalTensor。
• DataCopy() 把当前 Tile 从 GlobalTensor 搬入 LocalTensor。
• EnQue() 把数据放入输入 Queue。

Compute

• 典型过程：
• 从输入 Queue DeQue() 取出 xLocal、yLocal。
• 从输出 Queue AllocTensor() 获得 zLocal。
• 调用 Add(zLocal, xLocal, yLocal, count)。
• 将 zLocal EnQue() 到输出 Queue。
• 释放已经消费完的输入 LocalTensor。

CopyOut

• 典型过程：
• 从输出 Queue DeQue() 取得 zLocal。
• DataCopy() 将结果写回 Global Memory。
• FreeTensor() 释放输出 LocalTensor。

Double Buffer 机制

• Double Buffer 的核心是为流水阶段准备两组可交替使用的缓冲区。

• 没有 Double Buffer 时，一块 Local Memory 同时只能被一个阶段占用，CopyIn、Compute、CopyOut 容易互相等待。

• 采用双缓冲后：
• 当前 Tile 正在 Compute 时，可以为下一个 Tile 执行 CopyIn。
• 当前 Tile 正在 CopyOut 时，另一个 Tile 可以继续 Compute。
• 通过两个 Buffer 在不同 Stage 间轮换，降低硬件单元的空闲时间。

• 关键理解：Double Buffer 不是把数据“多算一遍”，而是用额外一份片上缓冲换取更高的流水并行度。

1-4 运行验证

• 一个完整的 Add Kernel 验证通常包括：
• 生成输入数据。
• 分配输入、输出内存。
• 调用核函数。
• 将结果写出。
• 用 CPU/NumPy 等参考实现计算 Golden 结果。
• 比较设备结果和 Golden 结果。

• Kernel 直调示例既可以运行 CPU 仿真，也可以运行 NPU。CPU 模式便于通过普通 C/C++ 调试方式定位逻辑问题；NPU 模式用于验证真实设备上的执行结果和性能行为。

2-Host侧实现

2-1 Host 侧实现概述

Host 与 Device

• 在典型昇腾算子运行环境中：
• Host：与 Device 相连的 x86/ARM 服务器，负责应用程序、运行时管理和 Host 侧算子逻辑。
• Device：安装昇腾 AI 处理器的设备，通过 PCIe 等接口与 Host 连接，提供 NPU 计算能力。

• Kernel 负责 Device 侧计算，而标准自定义算子还需要 Host 侧逻辑描述“怎么切、输出是什么、算子是什么”。

Host 侧三部分核心实现

• Host 侧主要包括：
1. Tiling 实现：根据输入 Shape、数据类型和硬件资源计算切分参数。
2. Shape 推导：根据输入 Tensor 描述和算子属性推导输出 Tensor 描述。
3. 算子原型注册：定义输入、输出、数据类型、Format，并把 Tiling、Shape 推导等函数注册到算子。

• 可以把三者理解为：
• 算子原型注册：这个算子“是什么”。
• Shape 推导：输出“长什么样”。
• Tiling：运行时“怎么切、怎么并行”。

2-2 Tiling 下发

为什么需要 Tiling

• Local Memory 容量有限，通常无法一次放下算子的全部输入、输出数据。因此要把大 Tensor 切成多个较小的数据块，再分给多个核、多个 Tile 处理。

• Tiling 的目标是确定：

• 使用多少个 AI Core。
• 每个核处理多少数据。
• 每个核内部再切成多少 Tile。
• 单个 Tile 的大小。
• 尾块如何处理。
• 需要多少额外 Workspace。

• Tiling 通常在 Host CPU 上执行，因为它主要进行 Shape、长度和硬件资源相关的标量计算，不适合占用 AI Core。

Tiling 结构体

• Tiling 结果需要从 Host 传到 Kernel，因此需要一种双方都能理解的数据结构。

• 简单 Add 示例可以只包含：

• 标准自定义算子工程中通常通过 Tiling 宏定义数据字段，例如：

• Host 侧实例化并填充该结构，随后写入 Tiling Buffer；Kernel 侧读取同一份数据。

Tiling 函数的工作流程

• Tiling 函数一般完成以下工作：
1. 从 TilingContext 获取输入 Tensor 的 Shape 等信息。
2. 计算总元素数。
3. 决定 blockDim，即启用多少个 AI Core。
4. 计算单核数据量、Tile 数和 Tile 大小。
5. 设置 Workspace 大小（如果算子需要）。
6. 把 TilingData 写入运行时提供的 Tiling Buffer。
7. 设置实际 TilingData 长度并返回成功。

Kernel 侧读取 Tiling 信息

• 标准核函数通常把 Tiling 信息放在参数列表尾部，例如：

• 易错点：标准工程中参数顺序通常按输入、输出、Workspace、Tiling 排列，不要随意调整。

固定 Shape 与动态 Shape

• 固定 Shape 时，很多切分参数可以编译期写成常量；动态 Shape 时，这些值必须由 Host 在运行时计算并下发。

• 动态 Shape 改造的核心可以概括成：静态常量 → 由 Tiling 下发的成员变量

• 例如原来 Kernel 中写死的 BLOCK_LENGTH、TILE_NUM，动态 Shape 下应从 tilingData 得到。

固定 Shape 与动态 Shape 工程文件差异

2-3 Shape 推导

Shape 推导的意义

• 神经网络可以看作一个有向无环计算图，每一个结点都是一个算子。当前算子的输出通常又会作为后续算子的输入，因此运行前需要尽可能确定每个输出 Tensor 的描述信息。

• Shape 推导根据：
• 输入 Tensor 的 Shape。
• 输入 Tensor 的数据类型、Format。
• 算子属性。
• 算子的数学语义。

• 推导输出 Tensor 的 Shape 等信息。

Shape 推导的作用

1. 参数校验：尽早发现输入不满足算子约束的问题。

2. 输出描述推导：确定输出 Shape、dtype、Format 等信息。

3. 静态内存规划：计算图构建阶段即可为已知 Tensor 分配内存，减少运行时动态分配开销。

• 对于逐元素 Add，如果输入 x、y Shape 已经满足广播/匹配要求，最简单场景下输出 Shape 与输入 Shape 相同。

• 概念代码：

2-4 原型注册

• 算子原型注册用于完整描述算子的“接口契约”，主要包括：
• 算子名称。
• 输入名称与数量。
• 输出名称与数量。
• 参数是否必选。
• 支持的数据类型。
• 支持的 Format。
• Shape 推导函数。
• Tiling 函数。
• 对应的 AI Core / SoC 配置。

• Add 示例可抽象为：

• Host 侧原型、Shape 推导与 Tiling 最终在同一个算子定义中关联起来，使框架能够完成算子发现、校验、编译与运行。

3-算子开发工程

3-1 算子开发工程概述

• Kernel 和 Host 都写完后，还需要解决四个工程问题：
• 怎样组织代码和目录？
• 怎样编译 Kernel 与 Host？
• 怎样生成可安装的算子包？
• 怎样部署到运行环境？

• Ascend C 提供两种典型开发工程。

Kernel 直调工程：快速流程

• 适合：
• 学习 Ascend C Kernel。
• 快速验证算法逻辑。
• 调试 Kernel。
• 不希望先编写完整 Host 侧标准算子工程。

• 特点：
• 文件少。
• 开发周期短。
• Kernel 开发完成后可以直接写 Host 测试程序调用。
• Tiling 可以简单处理，不依赖完整 CANN 标准算子注册流程。

自定义算子工程：标准流程

• 适合：
• 正式开发可部署算子。
• 需要框架调用。
• 需要 Host 侧原型注册、Shape 推导、Tiling。
• 需要生成算子安装包。
• 后续需要通过 AscendCL、PyTorch Adapter 等方式调用。

• 特点：
• 工程文件更多。
• 开发流程更完整。
• 可以通过工程脚本统一完成编译、打包和部署。

3-2 快速流程——Kernel 直调工程

Kernel 直调的核心思想

• Kernel 直调就是：Kernel 写完后，不先构建完整自定义算子包，而是直接写一个 Host 测试程序调用核函数。

• 最简工程只需要：

• 为了形成完整的自动验证流程，通常进一步加入：

CPU 与 NPU 两种验证路径

• Kernel 直调支持两类运行方式：
• CPU 调测：使用 CPU 仿真相关接口，例如 ICPU_RUN_KF，适合快速调试逻辑。
• NPU 调测：使用核函数调用语法或 AscendCL Runtime API，在真实 NPU 上执行。

• CPU 调测可以使用常规 C/C++ 手段，例如：
• gdb
• printf
• std::cout

• NPU 调测更关注：
• printf
• DumpTensor
• 真实硬件行为与执行结果

run.sh 的典型流程

• Kernel 直调的价值在于减少工程噪声，让开发者先集中验证 Kernel 本身。

3-3 标准流程——自定义算子工程

标准流程与快速流程对比

Add 算子分析

• Add 的数学表达式：

• 实现时需要明确：
• 输入：x、y。
• 输出：z。
• 示例 dtype：half。
• 示例 Shape：(8, 2048)。
• 示例 Format：ND。
• 搬运 API：DataCopy。
• 计算 API：Add。
• 内部数据对象：LocalTensor。
• 流水管理：Queue、EnQue、DeQue。

使用 msopgen 创建算子工程

• CANN 提供 msopgen 工具，可根据算子描述 JSON 自动生成标准工程骨架。

• 基本步骤：
1. 编写算子描述 JSON。
2. 使用 msopgen 生成工程。
3. 在生成的 Kernel 目录中完成设备侧实现。
4. 在 Host 目录中完成原型、Shape 推导和 Tiling。
5. 编译并生成安装包。

• 算子描述 JSON 重点包含：
• 算子名称。
• 输入列表。
• 输出列表。
• Shape/Format 约束。
• dtype 约束。

标准工程的主要目录

• 标准工程一般可以概括为：

• Kernel 与 Host 被分别编译，随后由工程脚本打包成自定义算子安装包。

Kernel 侧实现

• 标准工程中的 Kernel 仍然遵循：

• 与快速工程的主要区别不是计算逻辑本身，而是：
• 切分参数从 TilingData 获取。
• Kernel 与 Host 原型建立正式关联。
• 编译产物进入自定义算子包。

Host 侧 TilingData 定义

• 标准工程中可通过宏定义 TilingData：

• 随后注册 TilingData 类型，并由 Host 的 Tiling 函数填充。

工程编译

• 编译阶段主要完成：
• Kernel 源码编译。
• Host 源码编译。
• Tiling / 原型相关代码编译。
• 生成自定义算子安装包 .run。

• 常见 CMake 配置项包括：

配置项	作用
ASCEND_CANN_PACKAGE_PATH	指定 CANN 安装路径
CMAKE_BUILD_TYPE	Release / Debug 等构建类型
ENABLE_SOURCE_PACKAGE	是否生成源码相关包
ENABLE_BINARY_PACKAGE	是否生成二进制相关包
vendor_name	自定义算子厂商/包命名信息

• 工程通常通过 build.sh 完成构建，并在 build_out 等目录生成安装包。

算子包部署

• 编译完成后执行生成的 .run 安装包，将自定义算子部署到目标环境。

• 部署后的文件通常按功能分布到：
• Host 配置与算子原型相关目录。
• Kernel 二进制/源码相关目录。
• Tiling、配置、版本等辅助目录。

• 这样运行时和框架才能发现并加载自定义算子。

4-API通用解读

4-1 Ascend C 编程 API 概述

基础 API

• 基础 API 更贴近硬件能力，主要包括：
• 计算 API。
• 数据搬运 API。
• 内存管理 API。
• 任务同步 API。

• 基础 API 的输入输出通常围绕 GlobalTensor 与 LocalTensor 展开，灵活度高，适合自己设计数据通路和计算过程。

高阶 API

• 高阶 API 对常见复杂算法进行进一步封装，例如：
• Matmul。
• Softmax。
• Sinh 等数学函数。

• 使用高阶 API 可以减少重复开发，但通常需要根据接口要求准备临时空间、参数对象或初始化流程。

4-2 基础 API

基础 API 四大类

计算 API 的三种使用层次

整个 Tensor 参与计算

• 最简单的写法直接对整个 LocalTensor 计算，例如：

• 适合连续数据和简单逐元素计算。

Tensor 前 N 个元素参与计算

• 当只需要处理当前 Tile 中的部分有效数据时，可以通过 count 控制实际元素数。

Tensor 高维切分计算

• 更复杂的 Vector API 可以通过以下参数控制多轮计算：
• repeatTimes
• repeatStride
• blockStride
• mask

• 它们用于表达跨 block、跨 repeat 的规则访问方式。

repeatTimes

• Vector 计算单元一次 Repeat 可处理若干连续 block。典型向量场景中：
• 一个 block 为 32 Byte。
• 一次 Repeat 可覆盖 8 个 block，即 256 Byte。
• 数据超过一次 Repeat 的容量时，通过 repeatTimes 增加重复次数。
• repeatTimes 的有效范围存在上限，设计 Tile 时要避免单次调用超过接口限制。

• 例如 16 个 block 的连续数据，可分成 2 次 Repeat 处理。

repeatStride

• repeatStride 描述相邻 Repeat 起始位置之间的地址跨度。

• 常见用途：
• 连续计算：下一个 Repeat 紧接上一个 Repeat。
• 间隔计算：不同 Repeat 之间存在空洞。
• 重复计算：stride 为 0 时，可反复读取同一片数据。

blockStride

• blockStride 描述同一次 Repeat 内不同 block 之间的地址跨度。

• 因此要区分：

Mask：连续模式

• Mask 决定一次向量指令中哪些元素真正参与计算。

• 连续模式直接指定前多少个元素有效：
• 16 bit 数据单次可覆盖的元素数更多。
• 32 bit 数据单次可覆盖的元素数相应减少。

• 适合“前 N 个元素有效”的场景。

Mask：逐 bit 模式

• 逐 bit Mask 用位图描述元素是否参与计算：

• 适合非连续选取元素的场景。

• 理解重点：repeatTimes / repeatStride / blockStride / mask 共同描述一次 Vector API 如何在 Local Memory 中“走地址”。

数据搬运 API：DataCopy

• 常见搬运方向包括：
• GM → Vector 输入区。
• Vector 输出区 → GM。
• GM → Cube 输入相关位置。
• Cube 计算结果位置 → GM。
• 不同片上逻辑位置之间的搬运。

• 对于最常见的连续数据：

• 复杂二维/跨步搬运可以使用 DataCopyParams。

DataCopyParams

• 可以把二维搬运理解成：

内存管理 API

• TPipe 负责管理片上内存资源，TQue 表示某个流水阶段使用的 Queue。

• 典型流程：

EnQue / DeQue

• EnQue() 与 DeQue() 不只是普通容器操作，而是流水阶段间的同步机制。

• Producer 把准备好的 Tensor 放入 Queue，Consumer 在数据可用后取出，从而维持 Stage 之间正确的数据依赖关系。

4-3 高阶 API

高阶 API 的价值

• 高阶 API 封装了常见算法的复杂实现。以 Matmul 为例，开发者通常只需要：
1. 创建 Matmul 对象。
2. 初始化。
3. 设置左矩阵 A。
4. 设置右矩阵 B 和 Bias 等参数。
5. 执行矩阵乘。
6. 结束计算。

• 相比手工组织 Cube 指令、搬运和同步，高阶 API 可以显著减少代码量。

高阶 API 的临时空间

• 部分高阶 API 需要额外临时空间。例如 Sinh 接口可能需要 sharedTmpBuffer 保存中间结果。

• 接口一般提供两种使用方式：
• 开发者显式传入临时空间。
• 使用无需显式临时空间参数的重载，由接口内部处理。

• 为了更好控制片上内存，Host 可以先计算高阶 API 所需临时空间的最小值和最大值，再把选择结果作为 Tiling 参数下发。

• 概念流程：

• 在允许范围内增大临时空间，有时可以带来更好的计算性能，但也会占用更多片上内存，因此需要在性能与内存之间权衡。

5-算子的多种调用方式

5-1 算子调用概述

两类工程对应两类调用体系

• 不同开发方式支持的调用路径不同。
• Kernel 直调工程：重点是直接调用核函数，适合开发和快速验证。
• 自定义算子工程：完成正式编译部署后，可以通过 AscendCL、模型执行或 PyTorch Adapter 等方式调用。

常见调用方式对比

• 推荐开发顺序：
1. 先在 CPU 侧验证算法和基本逻辑。
2. 再用 Kernel 直调在 NPU 上验证。
3. 最后切换到正式工程和框架调用。

5-2 Kernel 直调

CPU 侧运行

• CPU 仿真模式的基本流程：
1. 为输入、输出申请共享内存。
2. 从文件读取输入数据。
3. 通过 CPU Kernel 调用宏执行核函数。
4. 把输出写入文件。
5. 释放共享内存。

• 常见接口：

接口	用途
GmAlloc(size)	为 CPU 调测创建共享内存
ICPU_RUN_KF(...)	在 CPU 仿真环境调用 Kernel
GmFree(ptr)	释放共享内存

NPU 侧运行

• NPU 模式通常需要：
1. 初始化 AscendCL。
2. 设置 Device。
3. 创建 Context / Stream。
4. 申请 Host / Device 内存。
5. 把输入从 Host 拷到 Device。
6. 拉起 Kernel。
7. 把输出从 Device 拷回 Host。
8. 同步 Stream。
9. 释放内存、Stream、Context、Device。

• 因为 Kernel 调用通常是异步的，所以读取结果前要确保对应 Stream 已经同步完成。

输入数据生成

• gen_data.py 通常负责：
• 用 NumPy 随机生成 input_x、input_y。
• 计算 golden = input_x + input_y。
• 将输入与 Golden 结果写成 .bin 文件。

• 这样 Host 测试程序只负责读取二进制数据并执行 Kernel。

结果校验

• verify_result.py 读取：
• NPU/CPU Kernel 实际输出。
• Golden 输出。

• 再使用容差比较，判断结果是否通过。

Kernel 直调也可以使用 Tiling

• Kernel 直调并不意味着不能使用 Tiling。

• 与标准工程相比，Kernel 直调更自由：
• Tiling 数据可以自己定义结构体。
• 不一定要使用标准 Host 侧 Tiling 注册流程。
• 调用 Kernel 前直接构造 Tiling 数据并传入即可。
• Kernel 参数也不一定必须写成统一的 GM_ADDR tiling 形式，只要调用端和 Kernel 端保持一致。

• 因此它很适合先验证复杂切分逻辑，再迁移到标准工程。

5-3 通过 Ascend CL 调用算子

AscendCL 单算子调用的两种方式

• 完成标准自定义算子开发和部署后，可以使用 AscendCL 执行单算子。

• 主要有两种方式：
1. 单算子 API 执行（aclnn）：直接调用生成的单算子 API。
2. 单算子模型执行：先把算子描述编译成离线模型，再用 aclopExecuteV2 等接口执行。

两种方式的工程要求

源码编译与二进制编译

二进制编译

• 会对 Kernel 实现进行编译，生成：
• 算子描述信息。
• 相关 JSON。
• 算子二进制文件。

• 适合直接调用已经编译好的单算子二进制。

源码编译

• 保留 Kernel 源码，不预先把 Kernel 固化为最终二进制。模型转换阶段可由 ATC 等工具结合目标环境进行编译。

• 适合算子随模型一起转换、编译和加载的场景。

单算子 API（aclnn）调用

• 自定义算子二进制编译部署后，会生成相应的单算子 API。API 常采用两段式：

第一阶段：GetWorkspaceSize

• 作用：
• 校验参数。
• 计算本次执行需要的 Workspace 大小。
• 创建/返回执行器 executor。

第二阶段：执行接口

• 作用：
• 根据 workspaceSize 准备 Device Workspace。
• 传入 executor 和 Stream。
• 异步执行算子。

• 典型调用顺序：

单算子 API 调用程序编译

• 调用程序需要在 CMake 中配置：
• 自定义算子包的 include 目录。
• 自定义算子包的 lib 目录。
• AscendCL 相关 include/lib。
• 自定义单算子 API 库。
• 运行时依赖库。

• 最终生成调用程序后，即可直接调用安装好的自定义算子。

单算子模型执行

• 另一种方式是把算子先转成离线模型。

• 流程：

• 关键接口包括：

• JSON 描述需要明确：
• 算子类型。
• 输入名称。
• 输入 Shape。
• 输入 dtype。
• 输入 Format。
• 输出描述。

• ATC 的核心参数包括：
• -singleop：单算子描述 JSON。
• -output：输出模型目录/前缀。
• -soc_version：目标 AI 处理器型号。

5-4 通过 PyTorch 调用算子

PyTorch 适配的总体思路

• PyTorch 在训练和推理中会调用大量算子。Ascend Extension for PyTorch 的 op-plugin 提供算子适配机制，把 PyTorch 算子映射到昇腾侧实现。

• 适配主要包含两部分：
1. 算子注册分发：在 YAML 中描述算子定义和后端映射关系。
2. 适配插件实现：编写 C++ 适配代码，在 PyTorch 接口与昇腾单算子 API 之间做参数转换和调用。

工程准备

• 示例工程通常包含：

• C++ 适配层主要完成：
• 接收 PyTorch Tensor。
• 创建输出 Tensor。
• 调用昇腾侧单算子 API。
• 返回 PyTorch Tensor。

算子注册分发配置

• op_plugin_functions.yaml 用来描述：
• 算子是否属于官方算子或自定义算子。
• func：算子的函数签名，包括名称、输入和返回值。
• impl_ns：指定后端实现所对应的命名空间/调用路径。

• 对于自定义算子，需要先在该配置中加入函数定义，再为其提供实际 C++ 实现。

op-plugin 编译部署

• 基本过程：
1. 获取 op-plugin 源码。
2. 合入自定义算子的 C++ 实现和 YAML 配置。
3. 重新编译 op-plugin。
4. 安装生成的 Python wheel 包。
5. 在 PyTorch 中调用并验证。

• 示例安装流程可概括为：

PyTorch 调用测试

• 测试脚本通常：
1. 创建 CPU/PyTorch 参考输入。
2. 把输入移动到 NPU。
3. 调用注册后的自定义算子，例如：
4. 把结果转回 CPU。
5. 与参考结果比较。

• 完整链路因此是：

6-非对齐尾块处理

问题背景

• 前面的 Add 示例使用 (8, 2048) 的 half Tensor，可以比较自然地按多个 AI Core 平均切分，且每块数据容易满足 32 Byte 对齐。

• 如果输入变成：
• 情况就不同了。

• half 占 2 Byte，因此真实输入数据大小为：

• Ascend C 的很多数据搬运和向量计算要求按 32 Byte block 组织数据：

• 1320 Byte 不是 32 Byte 的整数倍，需要向上对齐：

• 即从元素角度看，相当于：

• 其中：
• 660 个是真实数据。
• 12 个只是对齐后的尾部空间，不应该作为有效输出参与最终语义。

多核切分：不能再简单平均

• 42 个 block 要分配给 4 个 AI Core：

• 因此可以采用：
• 2 个“大核”：每核处理 11 block。
• 2 个“小核”：每核处理 10 block。

• 即：

• 这意味着 Kernel 中不能再假定：

• 而要根据 block_idx 判断当前属于大核还是小核，从 TilingData 中读取不同的数据长度。

单核内部还要受 UB 容量限制

• 即使已经把数据分到多个核，每个核一次能处理多少数据还受到 Local Memory / UB 大小限制。

• 因此需要再次把单核数据切成多个 Tile：

• 大多数 Tile 可以使用相同长度，最后一个 Tile 可能是尾块。

复杂 TilingData 的设计

• 为了让 Kernel 正确处理“大核/小核 + 普通 Tile/尾 Tile”，Host 侧需要下发更多信息。

• 典型字段包括：

字段	含义
smallCoreDataNum	小核总处理数据量
bigCoreDataNum	大核总处理数据量
finalSmallTileNum	小核最后一个 Tile 的有效数据量/相关切分参数
finalBigTileNum	大核最后一个 Tile 的有效数据量/相关切分参数
tileDataNum	普通 Tile 的数据量
smallTailDataNum	小核尾 Tile 的有效数据量
bigTailDataNum	大核尾 Tile 的有效数据量
tailBlockNum	尾部对齐处理所需 block 数等尾块信息

• 具体字段可以根据实现调整，核心目标是让 Kernel 能回答三个问题：
1. 我这个核要处理多少数据？
2. 我要循环多少个普通 Tile？
3. 最后一个 Tile 到底有多少有效元素？

Host 侧复杂 Tiling 的计算思路

1. 获取 Shape 和 dtype

• Host 从 TilingContext 获取输入 Shape 和数据类型，计算：

2. 按 32 Byte 对齐

3. 根据核数分配 block

• 假设使用 coreNum 个核：
• 前 remainder 个核多拿 1 block。
• 其余核拿 baseBlocks 个 block。

• 由此得到大核和小核的处理长度。

4. 根据 UB 大小计算 Tile

• 再根据可用 UB 大小、双缓冲数量和 dtype 计算普通 Tile 能放多少元素。

5. 计算尾 Tile

• 分别计算大核和小核：

• 这些值全部写入 TilingData 下发给 Kernel。

Kernel 侧处理逻辑

• Kernel 的 Init() 不再使用一套固定常量，而是：
1. 读取 TilingData。
2. 获取 block_idx。
3. 判断当前是大核还是小核。
4. 计算当前核的 GlobalTensor 起始偏移。
5. 确定本核普通 Tile 数与尾 Tile 长度。
6. 初始化 Queue/Buffer。

• 伪代码：

• Process() 中普通 Tile 可以复用原来的 CopyIn → Compute → CopyOut 逻辑；最后一个 Tail Tile 则需要使用实际有效元素数处理。

为什么要区分“搬运长度”和“有效计算长度”

• 尾块处理最容易混淆的地方是：
• 搬运可能需要满足 32 Byte 对齐。
• 真实数学计算只应该覆盖有效元素。

• 例如真实只剩 4 个 half：

• 为了满足搬运要求，可能仍然需要按一个 32 Byte block 处理内存，但计算时应通过 count、Mask 或额外逻辑确保只有真实元素参与最终语义。

• 核心原则：对齐是硬件访问要求，不等于对齐补出来的数据也属于 Tensor 的有效数据。

多数据类型算子实现

• 复杂 Shape 处理完后，还需要考虑另一类通用性问题：同一个算子支持多种 dtype。

Host 侧注册多种 dtype

• 算子描述 JSON 和 Host 侧原型注册需要声明支持的数据类型。例如同一输入可能支持：
• half
• float
• int32
• int8

• 输入、输出的数据类型约束要保持一致或按算子语义明确转换关系。

Kernel 侧模板化

• Kernel 可通过模板类型统一实现：

• 这样可以复用大部分搬运、切分和流水代码。

计算 API 不支持当前 dtype 时怎么办

• 即使 Host 注册了某个 dtype，也不代表所使用的 Vector API 一定直接支持它。

• 例如某个 Add 基础 API 不直接支持当前 int8 输入时，可以采用：

• 实现时可以使用 if constexpr 针对特定模板类型走专用分支：

• 因此，“算子支持某 dtype”需要同时满足两层条件：
1. Host 原型允许该 dtype。
2. Kernel 内部存在可执行的计算路径。

非对齐尾块处理的完整思路

• 最终可以把这一章归纳为一条通用切分链路：

• 这套方法不仅适用于 (1, 660) 的 Add，也适用于更一般的动态 Shape、非对齐输入和多数据类型 Vector 算子。