1.基本概念

1.1 Ascend C

什么是Ascend C？
Ascend c是CANN针对算子开发场景推出的编程语言，原生支持C和C 标准规范，最大化匹配用户开发习惯;通过多层接口抽象、自动并行计算、李生调试等关键技术，极大提高算子开发效率，助力AI开发者低成本完成算子开发和模型调优部署。
使用Ascend C开发自定义算子的优势:

• C/C 原语编程，最大化匹配用户的开发习惯
• 编程模型屏蔽硬件差异，编程范式提高开发效率
• 多层级API封装，从简单到灵活，兼顾易用与高效
• 李生调试，CPU侧模拟NPU侧的行为，可优先在CPU侧调试

1.2 CANN

CANN是华为针对AI场景推出的异构计算架构，本次活动主要聚焦的是其中的算子开发的部分。

1.3应用场景

将host == 服务器，Device就是华为的NPU。而一张NPU中有多个Aicore核心。

Ascend C能够为华为AI加速卡在大规模神经网络计算加速。

1.4 AIcore

首先，既然提到了AIcore,那具体AIcore能做什么呢？

Aicore支持核心计算，分别是：

• 标量（scalar)
• 向量（vector）
• 矩阵（cube）

以上图为例，AI Core中包含计算单元、存储单元、搬运单元等核心组件。

• 计算单元包括了三种基础计算资源：*Cube计算单元、Vector计算单元和Scalar计算单元。
• 存储单元即为AI Core的内部存储，统称为Local Memory，与此相对应，AI Core的外部存储称之为Global Memory。
• DMA搬运单元负责在Global Memory和Local Memory之间搬运数据。

且针对存在在不同区域中的数据类型，不论其原本的数据类型是什么（int,float…)

我们将用于存放AI Core中Local Memory（内部存储）的数据成为Local Tensor,

将用于存放AI Core中Gocal Memory（内部存储）的数据成为Gocal Tensor,

1.5 并行计算常见模型

并行计算常见模型有两种，SPMD（Single-Program Multiple-Data）数据并行 和 流水线并行。

前者SPMD将数据切分成不同部分，经多个进程处理，最好一同输出。

后者流水线同样是将数据切分，同时将进程的任务拆分成多个任务，全部数据如流水线操作一般，与SPMD不同，每个进程只会专注于一个任务的处理，会处理所有的数据分片。

2.编程模型与范式

上面介绍的主要是Ascend C中的一些基础概念。接下来主要介绍编程模型与范式。

编程模型主要由三个部分组成：

• 并行编程SPMD
• 核函数
• API

2.1 并行编程SPMD

Ascend C算子编程是SPMD（Single-Program Multiple-Data）编程，具体到Ascend C编程模型中的应用，是将需要处理的数据被拆分并同时在多个计算核心（类比于上文介绍中的多个进程）上运行，从而获取更高的性能。多个AI Core共享相同的指令代码，每个核上的运行实例唯一的区别是block_idx不同，每个核通过不同的block_idx来识别自己的身份。block的概念类似于进程的概念，block_idx就是标识进程唯一性的进程ID。编程中使用函数GetBlockldx()获取ID。并行计算过程的示意图如下图所示。

2.2 核函数

从SPMD模型可以得知，使用Ascend C进行编程时，我们编写一份算子实现代码，算子被调用时，将启动N个运行示例，在N个核上运行。本节将介绍算子实现的入口函数。

核函数（Kernel Function）是Ascend C算子设备侧实现的入口。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，当核函数被调用时，多个核都执行相同的核函数代码，具有相同的参数，并行执行。

Ascend C允许用户使用核函数这种C/C 函数的语法扩展来管理设备端的运行代码，用户在核函数中进行算子类对象的创建和其成员函数的调用，由此实现该算子的所有功能。核函数是主机端和设备端连接的桥梁。

//核函数的声明
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z);

这里其实可以看出，核函数的声明和普通C 函数声明大有不同。其中

global和aicore是函数类型限定符，使用global函数类型限定符来标识它是一个核函数，可以被<<<…>>>调用；使用aicore函数类型限定符来标识该核函数在设备端AI Core上执行。参数中的gm则表示存储在Global memory中。

编程中使用到的函数可以分为三类：核函数（device侧执行）、host侧执行函数、device侧执行函数（除核函数之外的）。三者的调用关系如下图所示：

• host侧执行函数可以调用同类的host执行函数，也就是通用C/C 编程中的函数调用；也可以通过<<<>>>调用核函数。
• device侧执行函数（除核函数之外的）可以调用调用同类的device执行函数。
• 核函数可以调用device侧执行函数（除核函数之外的）。

这里也可以看出核函数是作为host侧核Device侧之间的桥梁，让两边的执行函数连接起来。

除此之外，还有两条核函数应该遵守的规则：

1. 核函数必须具有void返回类型。
2. 仅支持入参为指针或C/C 内置数据类型（Primitive data types），如：half* s0，float* s1、int32_t c。

//调用核函数
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
// blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数，该核函数的参数列表为x，y，z。
add_custom<<<8, nullptr, stream>>>(x, y, z);

执行配置由3个参数决定：

• blockDim，规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即blockidx，可以在核函数的实现中调用[GetBlockIdx](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/70RC1alpha003/operatordevelopment/ascendcopdevg/atlasascendcapi070129.html)来获取block_idx；
• l2ctrl，保留参数，暂时设置为固定值nullptr，开发者无需关注；
• stream，类型为aclrtStream，stream是一个任务队列，应用程序通过stream来管理任务的并行。

下方是Add算子的例子：

// 实现核函数
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();
}
// 调用核函数
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);
}

2.3 API

Ascend C算子采用标准C 语法和一组类库API进行编程，类库API主要包含以下几种，您可以在核函数的实现中根据自己的需求选择合适的API：

• 计算类API，包括标量计算API、向量计算API、矩阵计算API，分别实现调用Scalar计算单元、Vector计算单元、Cube计算单元执行计算的功能。
• 数据搬运API，上述计算API基于Local Memory数据进行计算，所以数据需要先从Global Memory搬运至Local Memory，再使用计算接口完成计算，最后从Local Memory搬出至Global Memory。执行搬运过程的接口称之为数据搬移接口，比如DataCopy接口。
• 内存管理API，用于分配管理内存，比如AllocTensor、FreeTensor接口。
• 任务同步API，完成任务间的通信和同步，比如EnQue、DeQue接口。不同的API指令间有可能存在依赖关系，从AI Core内部并行计算架构抽象可知，不同的指令异步并行执行，为了保证不同指令队列间的指令按照正确的逻辑关系执行，需要向不同的组件发送同步指令。任务同步类API内部即完成这个发送同步指令的过程，开发者无需关注内部实现逻辑，使用简单的API接口即可完成。

Ascend C API的计算操作数都是Tensor类型：GlobalTensor和LocalTensor。

这里简单理解，就是在调用同样功能的API时，0级的计算性能会比其它等级的API更好。（0>1>2>3)

2.4 编程范式

在有了上述的核函数，API，并行编程等作为工具之后，编程范式描述了算子实现的固定流程。

把算子核内的处理程序，分成多个流水任务，通过队列（Queue）完成任务间**通信和同步，并通过统一的内存管理**模块（Pipe）管理任务间通信内存。流水编程范式应用了流水线并行计算方法。它提供了：

• 快速开发编程的固定步骤
• 统一代码框架的开发捷径
• 使用者总结出的开发经验
• 面向特定场景的编程思想
• 定制化的方法论开发体验

以下图为例解释流水任务，其可以看作两种并行计算常见方法的组合，即将数据切分后，将线程任务也切分，通过多线程快速处理切分的数据。

Ascend C分别针对Vector、Cube编程设计了不同的流水任务，

• Vector编程范式把算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn负责搬入操作，Compute负责矢量计算操作，CopyOut负责搬出操作。
• Cube编程范式把算子的实现流程分为5个基本任务：CopyIn，Split，Compute，Aggregate，CopyOut。CopyIn负责搬入操作，Split负责数据切分操作，Compute负责矩阵指令计算操作，Aggregate负责数据汇聚操作，CopyOut负责搬出操作。

上文中提到，进行编程范式需要将数据切分，不同的流水任务之间存在数据依赖，那如何保持任务间通信和同步？

Ascend C中使用Queue队列完成任务之间的数据通信和同步，提供EnQue、DeQue等基础API。我们以矢量（vector)编程中的流程为例。矢量编程中使用到的逻辑位置（QuePosition）定义如下：

• 搬入数据的存放位置：VECIN；
• 计算中间变量的位置：VECCALC；
• 搬出数据的存放位置：VECOUT。

1. Stage1：CopyIn任务。

• 使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。使用EnQue将LocalTensor放入VECIN的Queue中。

1. Stage2：Compute任务。

• 使用DeQue从VECIN中取出LocalTensor。使用Ascend C接口完成矢量计算。使用EnQue将计算结果LocalTensor放入到VECOUT的Queue中。

1. Stage3：CopyOut任务。

• 使用DeQue接口从VECOUT的Queue中取出LocalTensor。使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。

cube的编程范式与Vector类似，只是多了spilt和aggreagte的环节。

且对于VECIN 和 VECOUT 等queue的创建和删除，任务间数据传递使用到的内存统一由内存管理模块Pipe进行管理。如下图所示，Pipe作为片上内存管理者，通过InitBuffer接口对外提供Queue内存初始化功能，开发者可以通过该接口为指定的Queue分配内存。

Queue队列内存初始化完成后，需要使用内存时，通过调用AllocTensor来为LocalTensor分配内存，当创建的LocalTensor完成相关计算无需再使用时，再调用FreeTensor来回收LocalTensor的内存。

这里和C/C 的内存管理有相似的地方，即new 和 delete[] 需要成对出现。这样对资源的管理在Ascend C 的编程中会经常看见。

3.Helloworld实例

//在代码中，由于需要分别在CPU和NPU中调式，所以我们会使用__CCE_KT_TEST__来表示不同的调用程序。
#ifdef __CCE_KT_TEST__
// 用于CPU调试的调用程序
#else
// NPU侧运行算子的调用程序
#endif

宏是个很好用的工具，在之后我们还会接触到宏函数和其它的一些宏定义。