Ascend C算子开发（入门）笔记

type: Post
status: Published
date: 2026/08/10
slug: AscendC-Introduction
summary: Ascend C算子开发（入门）笔记，包含 算子基本概念；Tensor、Shape、Format 与 Axis；算子运行演示；算子开发的问题与挑战；CANN 与 Ascend C；昇腾 AI 处理器架构；Ascend C 的特点；Host 与 Device；核函数；Hello World 算子实现；完整核函数实现；算子开发环境部署概述；CPU 上部署开发环境；香橙派上部署开发及运行环境；华为云 ModelArts 上部署开发与运行环境
tags: Ascend, 学技术, 推荐, 文字, 机器学习
category: Ascend

@ZZHow(ZZHow1024)

参考课程：

【Ascend C算子开发（入门）】

[https://www.hiascend.com/developer/courses/detail/1691696509765107713]

【Ascend C系列教程（初级）】

[https://www.bilibili.com/video/BV1QgigeYEoz]

1-什么是算子

1-1 算子基本概念

人工智能的起点与基本层次

• 1956 年，美国达特茅斯人工智能研讨会召开，“人工智能（Artificial Intelligence）”概念由此正式进入研究视野。会议相关的代表人物包括约翰·麦卡锡（J. McCarthy）、马文·明斯基（M. L. Minsky）、克劳德·香农（C. E. Shannon）和 N. Rochester。
• 人工智能的目标，是让机器表现出类似人的智能行为，可以完成感知、认知、决策和执行等活动。按能力层次可以理解为：
  • 计算智能：能够进行计算，具备较强的存储能力和较快的计算能力。
  • 感知智能：能够感知视觉、听觉、触觉等环境信息，例如“会听、会说”。
  • 认知智能：能够理解、思考和决策。
  • 行动智能：能够自主学习、自主决策，并根据决策执行动作。

人工智能的三大流派

• 人工智能研究的发展过程中形成了符号主义、连接主义和行动主义三类典型思路。

流派	核心思想	代表方向/特点
符号主义	用符号、知识和逻辑规则描述认知过程	逻辑推理、启发式搜索、专家系统、知识工程
连接主义	用人工神经网络模拟神经系统的信息处理过程	BP 反向传播、卷积神经网络、循环神经网络、生成式神经网络等
行动主义	通过智能体的实际行为体现智能	机器人自主行动等场景

符号主义

符号主义认为，人类认知和思维的基本单元是“符号”，基于符号的一系列运算构成认知过程。计算机可以通过符号运算模拟人的智能活动。

• 优点：依赖知识推理，可解释性强。
• 缺点：过度依赖专家知识和规则，不具备从数据中学习的能力，可扩展性较弱。
• 代表成果：逻辑理论家、启发式搜索、专家系统、知识工程等。
• 典型案例：IBM Watson 自动问答系统。

连接主义

连接主义以人工神经网络研究为代表，通过构造人工神经网络模拟人脑的信息处理方式。它认为人的认知过程可以看作大量简单神经元构成的网络中的信息处理过程，而不只是符号运算。

• 代表成果：深度神经网络、BP 反向传播、卷积神经网络、循环神经网络、生成式神经网络等。
• 典型事件：2015 年 ImageNet 图像识别挑战赛中，机器在识别正确率上达到很高水平。
• 优点：能够从海量数据中学习，在图像处理、语音识别和自然语言处理等领域表现突出。
• 缺点：泛化能力有限、依赖大量样本，并且推理能力和可解释性相对不足。

符号主义与连接主义融合

符号主义擅长“推理”，连接主义擅长“学习”。将知识推理与数据学习结合，是人工智能发展的重要方向。

• 南京大学周志华教授提出“反绎学习（Abductive Learning）”，希望在一个框架内让机器学习与逻辑推理更均衡地协同。
• 清华大学张钹院士提出第三代 AI 的雏形，强调把数据驱动与知识推理结合，并进一步与人脑认知融合。

从生物神经元到人工神经元

生物神经元可以粗略理解为“树突接收信息 → 细胞核处理信息 → 轴突输出信息 → 突触传递信息”。人工神经元使用类似的抽象：

• 输入对应接收消息。
• 加权求和对应信息处理。
• 激活函数提供非线性变换。
• 输出继续传递给下一层神经元。

一个典型人工神经元的计算可以写成：

$$a=\sum _i{w}_i{x}_i+b$$

$$h=f(a)$$

其中，$x_i$ 是输入，$w_i$ 是权重，$b$ 是偏置，$f$ 是激活函数。人工神经元的核心可以概括为“线性组合 + 非线性激活”。

前馈神经网络

前馈神经网络中，各层神经元按照输入层、隐藏层、输出层组织：

• 第 0 层称为输入层。
• 最后一层称为输出层。
• 中间层称为隐藏层。
• 网络中不存在反馈回路，信号从输入层单向传播到输出层，因此可以用有向无环图表示。

隐藏层可表示为：

$$h^{(1)}=G\left((W^{(1)}{)}^{T}x+b^{(1)}\right)$$

$$h^{(l+1)}=G\left((W^{(l+1)}{)}^T{h}^{(l)}+b^{(l+1)}\right)$$

输出层可表示为：

$$\hat{y}=O\left((W^{(n)}{)}^T{h}^{(n)}+b^{(n)}\right)$$

神经网络的一个重要优势是特征表示学习：从原始、低层特征逐渐学习出中层和高层特征。例如在人脸相关任务中，可以从斑点、边缘逐渐形成鼻子、眼睛、脸颊等局部特征，最终形成“面部”这一更高层语义表示。

输出层与 Softmax

在分类任务中，输出层的作用可以看作分类器 $\hat{y}=f(x,w)$：

• $x$：最后一个隐藏层输出的高层特征。
• $w$：分类器参数。
• $\hat{y}$：各类别对应的输出概率。

对于 $K$ 分类问题，可以使用 Softmax 回归分类器：

$$p(y=k\mid x)=\mathrm{softmax}(w_k^{T}x)=\frac{\exp (w_k^{T}x)}{{\sum }_{k^{\prime }=1}^K\exp (w_{k^{\prime }}^{T}x)}$$

其中，每个类别 $k$ 对应一个参数向量 $w_k$。网络最后一层通常是包含 $K$ 个神经元的全连接层，再通过 Softmax 得到每个类别的条件概率。

参数学习

前馈神经网络中的参数主要包括各隐藏层和输出层的权重参数。

例如，一个四分类任务中，输入特征维度为 4，网络包含 3 个隐藏层，隐藏层神经元个数分别为 5、10、5，且暂不考虑偏置项，则参数数量为：

$$4\times 5+5\times 10+10\times 5+5\times 4=140$$

分类任务常使用交叉熵损失函数。对单个样本 $(x,y)$：

$$l(y,\hat{y})=-y^T\log \hat{y}=-\sum _{k=1}^K{y}_k\log {\hat{y}}_k$$

其中， y ∈ { 0 , 1 } K y\in\{0,1\}^{K} y∈{0,1}K 是标签的 one-hot 向量，$\hat{y}$ 是预测的类别概率向量。

整个训练集上的损失函数为：

$$L(W)=\frac{1}{N}\sum _{i=1}^{N}l\left(y^{(i)},{\hat{y}}^{(i)}\right)$$

梯度下降每次迭代时，第 $l$ 层参数的更新形式为：

$$W^{(l)}\leftarrow W^{(l)}-\alpha \frac{\partial L(W)}{\partial W^{(l)}}$$

反向传播与链式求导

反向传播的核心目标，是利用链式求导高效计算损失函数对网络参数的梯度。

如果：

$$y=g(x),z=f(y)=f(g(x))$$

标量情况下的链式法则为：

$$\frac{\partial z}{\partial x}=\frac{\partial z}{\partial y}\frac{\partial y}{\partial x}$$

扩展到向量形式，若 $x\in {\mathbb{R}}^m$、$y\in {\mathbb{R}}^n$：

$$\frac{\partial z}{\partial x_i}=\sum _j\frac{\partial z}{\partial y_j}\frac{\partial y_j}{\partial x_i}$$

写成向量形式：

$$\frac{\partial z}{\partial x}={\left(\frac{\partial y}{\partial x}\right)}^T\frac{\partial z}{\partial y}$$

其中 $\frac{\partial y}{\partial x}$ 是 Jacobian（雅可比）矩阵。

在神经网络中，可以把前向计算表示成计算图。反向传播时，从损失函数开始沿计算图反向传播梯度，逐层计算参数梯度。

自动微分

• 神经网络通常使用梯度下降优化参数。理论上可以手工按照链式法则逐个计算梯度，但手工求导并转换成程序过程繁琐、容易出错，且开发效率低。
• 主流深度学习框架因此提供自动微分能力：
  • 只需要描述网络结构和前向计算。
  • 框架根据计算图和链式求导规则自动计算梯度。
  • TensorFlow、PyTorch、MindSpore 等都支持自动微分，但具体实现方式可能存在差异。

算子在神经网络中的含义

• 在神经网络和计算图中，算子对应计算图中的一层或一个节点的计算逻辑。也可以理解为：每一个数据处理/数据计算的节点就是一个算子。

算子在神经网络计算图中的表现

• 一个神经网络通常由多个算子连接组成，数据 Tensor 在算子之间流动。算子接收输入数据，完成某种计算，再产生输出数据。

1-2 算子基本概念

算子的数学含义

• 数学意义上的算子可以理解为一个函数空间到函数空间的映射：

  $$O:X\to X$$

• 广义地说，对函数执行某种操作都可以看作一个算子，例如微分算子、不定积分算子等。在深度学习中，tanh、ReLU、sigmoid 等函数也可以对应为具体计算算子。

算子基本组成

算子名称（Name）

• 算子名称用于标识网络中的某个算子，同一网络中的算子名称需要保持唯一。
• 例如，一个网络中可以存在 Conv1、Pool1、Conv2。其中 Conv1 和 Conv2 名称不同，但都可以属于 Convolution 类型。

算子类型（Type）

• 算子类型决定算子的实现逻辑。相同类型的算子实现逻辑相同，同一个网络中可以存在多个同类型算子。

数据容器（Tensor）

• 算子执行时需要输入数据，执行完成后会产生输出数据。承载输入、输出数据的容器称为 Tensor（张量）。
• Tensor 是实际数据的容器，而 TensorDesc 是对输入、输出数据的描述。

Tensor 与 TensorDesc

属性	定义
名称（name）	用于对 Tensor 进行索引，不同 Tensor 的 name 需要保持唯一
形状（shape）	描述 Tensor 的形状，例如 (10,)、(1024, 1024)、(2, 3, 4) 等
数据类型（dtype）	指定 Tensor 对象的数据类型，例如 float16、float32、int8、int16、int32、uint8、uint16、bool 等；不同计算操作支持的数据类型可能不同
数据排布格式（format）	数据的物理排布格式，用于定义如何解释各个维度的数据

Shape

• 张量的形状通常写作：

  $$(D_0,D_1,\dots ,D_{n-1})$$

• 括号中有多少个维度值，就代表该 Tensor 是多少维。每一个维度的值表示该维度包含多少个元素。

  张量示例	Shape
  1	(0,)
  [1,2,3]	(3,)
  [[1,2],[3,4]]	(2,2)
  [[[1,2],[3,4]], [[5,6],[7,8]]]	(2,2,2)

• 例如 shape=(4,20,20,3) 可以理解为：

  • 有 4 张图片。
  • 每张图片的高度为 20、宽度为 20，即每张图有 $20\times 20=400$ 个像素。
  • 每个像素由 3 个通道组成，例如 RGB 三通道。

• 从程序角度看，多维 Tensor 可以理解为嵌套的多层循环；对某个元素的访问最终会映射为线性内存中的地址计算。

Format

• 深度学习中的多维数据最终仍需要在线性内存中存储，因此维度顺序会影响数据的物理排布。

• 卷积神经网络中的 Feature Map 常使用 4D 格式，其中：

  • N：Batch 数量，例如图片数量。
  • H：Height，特征图高度。
  • W：Width，特征图宽度。
  • C：Channels，特征图通道数，例如 RGB 图像的通道数为 3。

• 常见排布包括：

  Format	维度顺序	特点
  NCHW	[Batch, Channels, Height, Width]	同一通道的数据在内存中更集中，例如 RGB 图像可表现为先连续存 R，再连续存 G，再连续存 B
  NHWC	[Batch, Height, Width, Channels]	通道维位于最内层，同一像素位置的多个通道值相邻存储，例如 RGBRGBRGB...

axis

• axis 表示 Tensor 中某一个维度的下标。

• 如果 Tensor 的 shape=(5,6)：

  • axis=0 表示第一维，即“行”。
  • axis=1 表示第二维，即“列”。

• 例如数据：

  [[[1,2],[3,4]], [[5,6],[7,8]]]
  

• 其 shape=(2,2,2)：

  • 轴 0 对应最外层的两个矩阵。
  • 轴 1 对应 [1,2]、[3,4]、[5,6]、[7,8] 这一级数据。
  • 轴 2 对应最内层的标量 1,2,3,4,5,6,7,8。

• axis 也可以使用负数，从最后一个维度反向编号。对于 shape=(4,20,20,3)：

  正向 axis	负向 axis
  0	-4
  1	-3
  2	-2
  3	-1

• 对于 N 维 Tensor，正向轴编号为 0,1,2,...,N-1。

1-3 算子运行演示

• run.sh 的执行过程可以概括为：
  1. 编译 add_custom 算子。
  2. 编译算子调用方法 main.cpp。
  3. 执行 scripts 目录中的数据生成逻辑，调用 main.cpp 拉起算子计算，并将算子结果与预期结果进行比较，判断计算是否正确。
• 因此，一个最小算子样例通常包含“算子实现 → 编译 → Host 侧调用 → 输入数据准备 → 执行 → 结果校验”这一完整链路。

1-4 算子开发的问题与挑战

算子开发的复杂性

• 实现一个算子时，开发者不仅要“把数学公式写成代码”，还要同时考虑多个层面：
  • 功能逻辑如何实现。
  • 如何处理不同大小的输入。
  • 如何处理不同类型的输入。
  • 如何适配目标硬件。
  • 如何保证算子运行性能。
  • 如何优化算子的数学公式与计算过程。

功能逻辑实现：以激活函数为例

• 即使都是“激活函数”，不同函数的适用场景、性能特征以及对精度的影响也并不相同。因此，算子设计不能只关注“能否算对”，还需要考虑同一功能可能存在的多种实现方式。
• 典型激活函数包括：
  • tanh
  • ReLU
  • sigmoid

与硬件结合：以 Flash Attention 为例

• 高性能算子需要结合硬件存储层次和计算资源设计。Flash Attention 的核心思路之一，是减少高层、低带宽存储与片上高速存储之间的数据搬运，把更适合局部计算的数据块放到更靠近计算单元的高速存储中处理。
• 在昇腾硬件上，还需要同时考虑：
  • 核间并行：把不同数据块分配给多个 AI Core 并行执行。
  • 核内并行：协调 Scalar、Vector、Cube、DMA/搬运等不同单元，使计算与数据搬运尽可能重叠。

Flash Attention 与硬件存储/计算资源的结合

性能优化：提高流水并行度

• 以 Flash Attention 的性能优化为例，在对缓存中的 mm1/mm2/mm3 等计算进行优化后，可以在本轮 Vector 与 Cube 流水的间隔中，提前插入下一轮循环的 Vector 计算。
• 这样可以让 Vector 流水和 Cube 流水之间的并行度更高，在流水图中表现为 Vector 计算更加密集，从而提升整体资源利用率和执行性能。

2-什么是Ascend C

2-1CANN 与 Ascend C

CANN 在昇腾 AI 软件栈中的位置

• CANN 是面向昇腾 AI 处理器的异构计算架构，向上承接深度学习框架、AI 框架适配、创新算子与领域加速库、人工智能应用，向下连接 Runtime、驱动和昇腾 AI 处理器。
• 在 CANN 中，典型组件包括：
  • GE 图引擎：计算图编译运行控制中心，提供图编译优化与加载执行能力。
  • Ascend C 算子开发语言：面向算子开发场景，支持算子级编程。
  • AOL 算子加速库：提供经过深度优化的高性能算子。
  • HCCL 集合通信库：提供单机多卡及多机多卡的数据并行、模型并行等集合通信能力。
  • Runtime 运行时：提供资源管理、媒体数据预处理、模型推理等基础能力。
  • MindStudio：提供全流程开发工具链。

CANN 与 Ascend C 在昇腾 AI 软件栈中的位置

什么是 Ascend C

• Ascend C 是 CANN 针对算子开发场景推出的编程语言。它通过多层接口抽象、自动并行计算、孪生调试等关键技术，提高算子开发效率，帮助开发者以较低成本完成算子开发和模型调优部署。
• 使用 Ascend C 编程语言开发的算子称为 Ascend C 算子。

使用 Ascend C 开发自定义算子的优势

• C/C++ 原语编程：最大化匹配开发者已有的 C/C++ 开发习惯。
• 屏蔽硬件差异：通过编程模型抽象硬件差异，提高开发效率。
• 多层级 API 封装：从灵活的底层控制到高层易用接口，兼顾灵活性与效率。
• 孪生调试：可在 CPU 侧模拟 NPU 侧行为，优先在 CPU 环境中进行功能和精度调试。

2-2 昇腾 AI 处理器架构

昇腾 AI 处理器逻辑架构

• 昇腾 AI 处理器的逻辑架构主要包括：
  • 系统控制处理器（Control CPU）：负责系统控制相关任务。
  • AI Core：面向计算密集型任务的 AI 计算核心。
  • AI CPU：面向非矩阵类计算任务的 AI 处理器。
  • 任务调度器 TS：负责相关任务调度。
  • 层次化片上缓存/缓冲区：为计算单元提供高带宽数据访问。
  • 数字视觉预处理模块（DVPP）：完成数字视觉相关预处理。
  • I/O 接口：负责外部数据与设备连接。
  • DDR/HBM 接口：连接外部高容量内存。

昇腾 AI 处理器逻辑架构

AI Core 与达芬奇架构

• AI Core 是昇腾 AI 处理器的计算核心，采用华为自研的达芬奇架构（DaVinci Core）。不同处理器版本中的计算、存储和带宽资源规格可能不同，但总体可以划分为三大部分：
  • 计算单元：包含矩阵计算单元、向量计算单元、标量计算单元等基础计算资源。
  • 存储系统：由 AI Core 的片上存储单元以及相应数据通路组成，为计算单元提供数据。
  • 控制单元：负责整个计算过程的指令控制，相当于 AI Core 的“司令部”。
• 从具体数据通路看，AI Core 中会涉及 L1 Buffer、Unified Buffer、矩阵输入/输出 Buffer、Vector/Scalar 计算资源以及搬运单元等。

耦合架构与分离架构

• Ascend AI Core 可以采用不同的计算资源组织方式。

耦合架构

• Cube、Vector、Scalar 等资源集成在同一 AI Core 中，数据可通过 L1 Buffer、Unified Buffer 等片上存储协同流动。

分离架构

• 计算资源可进一步拆分为不同类型的核心：
  • AIC：以 Cube 矩阵计算为主要资源。
  • AIV：以 Vector 向量计算为主要资源。
• AIC 与 AIV 可按照一定比例组织，使矩阵计算和向量计算资源更加独立地调度和利用。

2-3Ascend C 的特点

开发效率提升

• 传统算子开发存在较高门槛，典型难点包括：
  • 程序语义如何映射到复杂指令序列。
  • 数据存储空间如何分配、释放和复用。
  • 如何实现数据和计算流水并行。
• Ascend C 通过编程抽象降低这些门槛，典型场景中可以显著缩短算子开发周期。
• 其核心特点可以概括为：
  • 遵循 C/C++ 标准规范。
  • 自动化流水并行调度。
  • 结构化核函数编程。
  • CPU/NPU 孪生调试。

采用标准 C++ 语法，基于类库 API 编程

• Ascend C 使用标准 C++ 语法，并通过类库 API 提供算子开发能力。常见 API 类型包括：

  • 向量计算 API。
  • 矩阵计算 API。
  • 数据搬运 API。
  • 内存管理 API。
  • 任务同步 API。

• 基本数据类型包括 GlobalTensor、LocalTensor 等。

• 计算 API 采用分层设计：

  API 层级	主要特点	示例
  0 级 API	功能灵活，可显式控制较多底层操作	Add(dst, src1, src2, mask, repeatTimes, repeatParams)
  1 级 slice 计算 API	面向多维数据切片计算	解决多维数据切片计算问题
  2 级连续计算 API	对 Tensor 指定长度的连续数据进行计算	Add(dst, src1, src2, count)
  3 级 API	运算符重载，表达更接近普通 C++	dst = src1 + src2

• API 层级越高，一般自由度越低，但易用性越高；层级越低，控制粒度更细、自由度更高。

核间支持 SPMD 数据并行

• Ascend C 支持 **SPMD（Single-Program Multiple-Data）**数据并行：
  • 将待处理数据拆分并分发到多个计算核心。
  • 多个 AI Core 运行相同的指令代码，但处理不同的数据分片。
  • 每个核通过不同的 block_idx 区分自己需要处理的数据。
  • 开发者重点关注单核算子实现，再由并行模型将工作扩展到多核。

核内支持自动化流水并行

• Ascend C 将算子核内处理过程拆分为多个流水任务（Stage），典型阶段是：
  1. 搬入（CopyIn）：将输入数据从 Global Memory 搬入 Local Memory。
  2. 计算（Compute）：使用 Local Memory 中的数据进行计算。
  3. 搬出（CopyOut）：将计算结果从 Local Memory 搬回 Global Memory。
• 其中：
  • Tensor 作为数据载体。
  • Queue 用于不同任务之间的通信和同步。
  • Pipe 用于管理任务间的通信内存。
• 通过多片数据流水调度，可以让“搬入、计算、搬出”阶段交叠执行，提高计算与搬运资源的并行度。

Ascend C 核内流水并行

结构化核函数编程

• Ascend C 提供结构化的算子实现框架，典型逻辑为：

  kernel_name
  ├── Init
  └── Process
      ├── CopyIn
      ├── Compute
      └── CopyOut
  

• 各阶段职责：

  • Init：完成内存初始化、队列等资源创建。
  • CopyIn：输入数据从 Global Memory 搬到 Local Memory。
  • Compute：使用 Local Memory 中的数据完成计算。
  • CopyOut：结果从 Local Memory 搬回 Global Memory。

• 这种结构有利于快速搭建算子实现代码框架，并与流水并行模型自然对应。

异构混合编程：Host/Device 灵活通信

• 昇腾 AI 异构计算中，CPU 与 NPU 协同工作。Ascend C 编程模型通常需要分别实现：

  • Host 侧代码：负责初始化、资源申请、数据传输、核函数启动、同步、资源释放等。
  • Device 侧代码：在 AI Core 上执行具体算子计算。

• 典型执行流程：

  1. AscendCL 初始化。
  2. 申请运行管理资源。
  3. Host 数据传输到 Device。
  4. 调用核函数完成指定运算。
  5. Device 数据传回 Host。
  6. 释放运行管理资源。
  7. AscendCL 去初始化。

• 核函数通过扩展调用语法启动：

  kernel_name<<<blockDim, nullptr, stream>>>(argument_list);
  

• 其中：

  • blockDim：执行核数。
  • nullptr：保留参数。
  • stream：用于异步任务调度的任务流。

CPU/NPU 孪生调试

• 传统方式在 NPU 环境中调试时，常面临调试耗时长、并发和地址问题定位困难等问题。Ascend C 支持 CPU/NPU 孪生调试：
  • CPU 域调试：重点验证功能和精度，可使用 GDB、printf/cout、ASSERT 等方式，定位逻辑错误、数据计算错误和内存问题。
  • NPU 域调试：重点验证真实性能和硬件执行行为，可使用 Profiling 流水图、指令日志、数据日志、板上执行时间统计等手段，定位性能问题和算子同步问题。

3-算子开发初体验

3-1Host 与 Device

• 在 Ascend C 算子开发中，Host 与 Device 构成典型的异构计算系统。
  • Host：与 Device 相连接的 X86 或 ARM 服务器，负责运行 Host 侧程序，并利用 Device 提供的神经网络计算能力完成任务。
  • Device：安装了昇腾 AI 处理器的硬件设备，通过 PCIe 等接口与 Host 连接，提供 NPU 计算能力。
• Device 上包含一个或多个 AI Core，并通过 Global Memory（例如 DDR）存储大容量数据。

Host 与 Device 的关系

3-2 核函数

什么是核函数

• **核函数（Kernel Function）**是 Ascend C 算子 Device 侧的入口。Ascend C 通过扩展 C/C++ 函数语法来管理设备侧运行代码，开发者在核函数中实现算子逻辑，例如定义算子类及其成员函数来完成数据搬运与计算。

• 核函数是 Host 侧和 Device 侧之间的重要桥梁。

• 典型形式：

  __global__ __aicore__ void kernel_name(argument_list);
  

• 与 CUDA 的典型写法相比：

  // Ascend C
  __global__ __aicore__ void kernel_name(argument_list);
  
  // CUDA
  __global__ void kernel_name(argument_list);
  

• 核函数直接在设备侧执行。SPMD 编程模型允许核函数启动后，由多个计算核心并行执行同一份代码并处理不同数据。

函数类型限定符

函数类型限定符	执行位置	调用方式	备注
__global__	设备侧	由 <<<...>>> 调用	必须为 void 返回类型
__aicore__	设备侧	仅从设备端调用	表示在 AI Core 上执行

• 因此，Ascend C 核函数通常同时使用 __global__ 和 __aicore__：

  __global__ __aicore__ void kernel_name(argument_list);
  

变量类型限定符

• 为了统一指针入参类型，可以使用：

  __gm__ uint8_t*
  

• __gm__ 表示该指针变量指向 Global Memory 中的某一内存地址。

  变量类型限定符	内存空间	含义
  __gm__	Global Memory	表明该指针变量指向 Global Memory 上某处内存地址

• 常用宏：

  #define GM_ADDR__gm__uint8_t*__restrict__
  

• 核函数入参规则/建议：

  1. 核函数必须具有 void 返回类型。
  2. 入参支持指针类型或 C/C++ 内置基础数据类型，例如 half*、float*、int32_t 等。
  3. 可以使用 GM_ADDR 统一封装 Global Memory 指针，避免函数入参列表过长。

如何调用核函数

• 普通 C/C++ 函数调用：

  function_name(argument_list);
  

• Ascend C 核函数使用内核调用符：

  kernel_name<<<blockDim, nullptr, stream>>>(argument_list);
  

• 其中：

  • blockDim：规定核函数在多少个核上执行。每个核会被分配一个逻辑 ID block_idx，编号从 0 开始；算子实现中可使用 GetBlockIdx() 获取当前逻辑核 ID。
  • stream：类型为 aclrtStream，表示任务队列，应用程序通过 stream 管理任务并行。

• 核函数调用是异步、非阻塞的。Host 侧发起核函数后不会默认等待其立即结束，需要在需要结果或保证执行完成的位置显式同步，例如调用 aclrtSynchronizeStream(stream)。

  内核调用符 <<<...>>> 只在 NPU 模式编译时使用；CPU 模式下不能直接识别这一调用符号。

3-3Hello World

Device 侧核函数实现

• 一个最小的 Hello World 核函数可以写成：

  #include"kernel_operator.h"
  using namespace AscendC;
  
  extern "C" __global__ __aicore__ void hello_world()
  {
      PRINTF("Hello World!!!\n");
  }
  
  void hello_world_do(uint32_t blockDim, void* stream)
  {
      hello_world<<<blockDim, nullptr, stream>>>();
  }
  

• 这段代码包含三个关键部分：

  • 核函数定义：hello_world 通过 __global__ __aicore__ 声明为 AI Core 核函数。
  • 核函数实现：使用 PRINTF 在核函数中输出信息。
  • 核函数调用封装：hello_world_do 使用 <<<blockDim, nullptr, stream>>> 拉起核函数。

Host 侧调用流程

• Host 侧需要先完成 AscendCL 初始化和运行资源创建，再调用核函数：

  #include"acl/acl.h"
  
  extern void hello_world_do(uint32_t coreDim, void* stream);
  
  int32_t main(int argc, char const *argv[])
  {
      aclInit(nullptr);
  
      aclrtContext context;
      int32_t deviceId = 0;
      aclrtSetDevice(deviceId);
      aclrtCreateContext(&context, deviceId);
  
      aclrtStream stream = nullptr;
      aclrtCreateStream(&stream);
  
      constexpr uint32_t blockDim = 8;
      hello_world_do(blockDim, stream);
      aclrtSynchronizeStream(stream);
  
      aclrtDestroyStream(stream);
      aclrtDestroyContext(context);
      aclrtResetDevice(deviceId);
      aclFinalize();
      return 0;
  }
  

• 常见 AscendCL 接口及其作用：

  接口	作用
  aclInit	AscendCL 初始化
  aclrtSetDevice	指定目标 Device
  aclrtCreateContext	创建运行上下文
  aclrtCreateStream	创建任务流
  aclrtMallocHost / aclrtMalloc	申请 Host/Device 内存
  aclrtMemcpy	Host 与 Device 之间进行数据传输
  <<<...>>>	启动 Device 核函数
  aclrtSynchronizeStream	等待指定 stream 中任务执行完成
  aclrtFree / aclrtFreeHost	释放 Device/Host 内存
  aclrtDestroyStream / aclrtDestroyContext	释放运行管理资源
  aclrtResetDevice	复位 Device
  aclFinalize	AscendCL 去初始化

• 完整调用链可以记为：

  AscendCL 初始化
      ↓
  创建 Device / Context / Stream
      ↓
  准备并传输数据
      ↓
  启动核函数
      ↓
  Stream 同步等待
      ↓
  取回结果
      ↓
  释放资源
      ↓
  AscendCL 去初始化
  

3-4 完整核函数范讲

• 一个完整的向量计算核函数通常不会把所有逻辑堆在一个函数中，而是按照 Ascend C 流水编程范式拆分为多个阶段。

典型代码结构

Kernel 类
├── Init
├── Process
├── CopyIn
├── Compute
└── CopyOut

• 各函数职责：
  • Init：完成 GlobalTensor 地址绑定、Pipe/Queue/Buffer 等资源初始化。
  • Process：组织循环，按 Tile 依次调度 CopyIn → Compute → CopyOut。
  • CopyIn：从 Global Memory 获取输入数据，放入本地 Tensor，并通过队列入队。
  • Compute：从输入队列出队，在 Local Memory 上执行向量/矩阵计算，结果放入输出队列。
  • CopyOut：从输出队列出队，将结果写回 Global Memory，并释放本地 Tensor。

TPIPE 流水编程模型

• Ascend C 的 TPIPE 流水编程以“搬入、计算、搬出”为三个主要 Stage：

  Global Memory
      ↓ CopyIn
  Local Memory
      ↓ Compute
  Local Memory
      ↓ CopyOut
  Global Memory
  

• 不同 Tile 可以在流水线上交错运行，例如当第 0 块数据正在 Compute 时，第 1 块数据可以 CopyIn，而前一块数据可以继续 CopyOut。这种重叠执行能提高搬运单元与计算单元的利用率。

完整向量核函数中的 CopyIn/Compute/CopyOut 流水

4-算子开发环境搭建

4-1 环境部署概述

开发环境与运行环境

• 开发环境主要用于代码开发、编译、调测等活动，存在两类常见场景：
  • 场景一：在非昇腾 AI 设备上安装开发环境。
    • 可用于代码开发、编译等不依赖昇腾设备的活动。
    • 例如 ATC 模型转换、算子和推理应用程序的纯代码开发。
  • 场景二：在昇腾 AI 设备上安装开发环境。
    • 支持代码开发和编译。
    • 同时可以运行应用程序，或进行训练脚本的迁移、开发与调试。
• 运行环境部署在昇腾 AI 设备上，用于运行开发完成的应用程序，或者进行训练脚本的迁移、开发与调试。

CANN 开发环境与运行环境部署流程

CANN 相关安装包

• 在昇腾社区中可以从“产品 → CANN”获取相关安装包。社区版更新频率较高，适合开发者使用。

  软件包名称（示例版本）	说明
  Ascend-cann-nnrt_8.0.RC2.alpha002_linux-x86_64.run	x86 平台推理引擎软件包，适用于命令行方式安装场景
  Ascend-cann-amct_8.0.RC2.alpha002_linux-aarch64.tar.gz	ARM 平台模型小型化工具，适用于命令行方式安装场景
  Ascend-cann-amct_8.0.RC2.alpha002_linux-x86_64.tar.gz	x86 平台模型小型化工具，适用于命令行方式安装场景
  Ascend-cann-communitysdk_8.0.RC2.alpha002_linux-aarch64.run	CANN 社区算子开发工具包，适用于命令行安装场景
  Ascend-cann-communitysdk_8.0.RC2.alpha002_linux-x86_64.run	CANN 社区算子开发工具包，适用于命令行安装场景
  Ascend-cann-robotmiddleware_8.0.RC2.alpha002_linux-aarch64.run	机器人应用开发中间件 OpenHiva 软件包，适用于命令行安装场景
  Ascend-cann-toolkit_8.0.RC2.alpha002_linux-aarch64.run	ARM 平台开发套件软件包，适用于命令行安装场景
  Ascend-cann-toolkit_8.0.RC2.alpha002_linux-x86_64.run	x86 平台开发套件软件包，适用于命令行安装场景
  Ascend-cann-kernels-310p_8.0.RC2.alpha002_linux.run	CANN 算子二进制安装包，适用于命令行安装场景
  Ascend-cann-kernels-310b_8.0.RC2.alpha002_linux.run	CANN 算子二进制安装包，适用于命令行安装场景
  Ascend-cann-kernels-910_8.0.RC2.alpha002_linux.run	CANN 算子二进制安装包，适用于命令行安装场景
  Ascend-cann-kernels-910b_8.0.RC2.alpha002_linux.run	CANN 算子二进制安装包，适用于命令行安装场景
  Ascend-cann-nnae_8.0.RC2.alpha002_linux-aarch64.run	ARM 平台深度学习引擎软件包，适用于命令行方式安装场景
  Ascend-cann-nnae_8.0.RC2.alpha002_linux-x86_64.run	x86 平台深度学习引擎软件包，适用于命令行方式安装场景
  Ascend-cann-nnrt_8.0.RC2.alpha002_linux-aarch64.run	ARM 平台推理引擎软件包，适用于命令行方式安装场景

4.2 在 CPU 上部署开发环境

选择安装场景-软件安装-发行与安装-CANN9.1.0开发文档-昇腾社区

4.3 在香橙派上部署开发及运行环境

选择安装场景-软件安装-发行与安装-CANN9.1.0开发文档-昇腾社区

4.4 在华为云 ModelArts 上部署开发&运行环境

CANNLab - 我的环境 - 开源代码托管,代码协作 - AtomGit