【2024CANN训练营】Ascend C算子开发课程笔记

时间：2024年7月10日 地点：华为北京研究所
课程内容：Ascend C算子开发。

什么是算子（Operator）？

如果把模型比喻成为一幢建筑，那么算子就是它的砖块。成功搭起建筑不能够从烧砖开始，模型的开发也应当使用已有的算子来进行搭建。具体来举例，基础的算子有Add，复杂的有针对特定场景应用的融合算子。

关注算子，就还应当关注张量（Tensor）。算子的输入和输出即为张量，张量是一个多维的数组，是标量、向量、矩阵的高维扩展。从编程角度来讲，张量这一结构除了数据本身外，还包含了对数据属性的描述。需要关注的属性有：

• 数据类型（Type）：例如FP16、FP32、int4、int8等
• 形状（Shape）：张量的维度数以及各个维度的大小
• 数据排布的格式：数据排布各个维度的意义，例如图像数据中的N（batch size）、H（height）、W（width）、C（channel num）。
• 轴（Axis）：选取不同的轴可以得到不同shape。例如[[[1,2][3,4]],[[5,6],[7,8]]]中，
  取axis = 0 ，则得到两个矩阵 [[1,2][3,4]]和[[5,6],[7,8]]；取axis=1，则得到四个向量，即[1,2]、[3,4]、[5,6]和[7,8]；取axis=0，则得到8个数字，即1、2、3、4、5、6、7和8。

什么是Ascend C？

Ascend C是一种编程语言，偏重于算子开发场景的应用。它原生支持C/C++，对于有算子开发需求的开发人员而言十分友好。Ascend C采用的关键技术有：

• 多层接口抽象：简单灵活，屏蔽不同硬件平台之间的差异
• 自动并行计算：自动利用并行执行提升计算效率
• 孪生调试：在CPU侧模拟调试NPU侧行为

AscendC算子编程是SPMD（单程序多数据）的编程，利用自动流水并行调度提升算子性能。使用结构化函数编程即可，按照规定的编程范式即可进行算子的开发。

具体来讲，Ascend C是应用在CANN计算架构中的计算语言。那么什么是CANN呢？

什么是CANN？

CANN是华为提出的AI异构计算架构，包含多种模块，具有从上至下的完整技术栈，可激发出昇腾AI处理器的计算能力。CANN提供多层次的编程接口，可以支持用户构建__基于昇腾平台__的AI应用和业务。

CANN对开发人员提供API，基于NPU进行训练，可以使能CPU和NPU协同进行编码（专人专事，NPU负责矢量、矩阵运算，其他用CPU）。
Ascend C编程语言开发的算子搭建成模型，使用对应编译器，最终运行在AICORE上。AICORE简单来讲包含三种单元：计算单元、存储单元、搬运单元。计算单元进行标量、向量和矩阵的运算；搬运单元进行global memory和local memory之间数据的搬运；存储单元存储有模型所用的数据，分为内部存储和外部存储。本次课程主要关注于CANN框架中的AscendC编程语言和算子库部分。

Ascend C编程入门

在具体接触代码前，应当了解一个内容AscendC算子编程是SPMD的编程，在计算时，将任务拆分于不同的进程，使用block idx来对数据和运算进行标注。SPMD即为单进程多数据（如下图），启动一组进程（都运行相同的代码），而后将待处理的数据进行切分，把不同的数据分发给不同进程处理，每个进程都对自己的数据分片进行所有任务的处理。

此外，我们还应当对算子的工作原理进行了解（如下图）。

算子的计算分为三个步骤：首先从数据中经过CopyIn一个Tensor得到输入用算子，而后针对算子功能和特性进行计算，最终计算结果CopyOut到结果Tensor中去。这三个步骤，通过队列这一结构来实现不同阶段之间的同步，保证计算的正确性。

算子实现相关函数的调用关系如下：

host侧

host侧应当关注的一个结构体为Tiling。该结构体中记录的参数用于表示如何对输入的数据进行切分，决定了计算过程的一些细节。该结构体在host端实例化，通过指针传入kernel函数中，用于指导计算过程。（一般在Init中实现对tiling相关的计算，应当尽量减少kernel函数中多余的计算）

tiliing官方介绍：
大多数情况下，Local Memory的存储，无法完整的容纳算子的输入与输出，需要每次搬运一部分输入进行计算然后搬出，再搬运下一部分输入进行计算，直到得到完整的最终结果，这个数据切分、分块计算的过程称之为Tiling。根据算子的shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循环多少次）的计算程序，称之为Tiling实现。

workspace 也在host端定义，可用于算子之间的数据传递，一般在算子融合场景下应用内较多。相对于写死而言，tiling可以实现动态shape（动态划分计算过程）。

写死（直接初始化）：

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // seperate to 2 parts, due to double buffer

动态shape：

namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    ReduceMaxTilingData tiling;
    // 从attr获取reduceDim属性值，因为reduceDim是第一个属性，所以GetAttrPointer传入的索引值为0
    const gert::RuntimeAttrs* attrs = context->GetAttrs();
    const uint32_t* reduceDim = attrs->GetAttrPointer<uint32_t>(0);
    // 获取reduceDim轴的长度
    const gert::StorageShape* xShapePtr = context->GetInputShape(0);
    const gert::Shape& xShape = xShapePtr->GetStorageShape();
    const uint32_t reduceAxisLen = xShape.GetDim(*reduceDim);
    // 计算TilingData中除了reduceAxisLen之外其他成员变量的值
    ...
    // 将reduceAxisLen设置到tiling结构体中，传递到kernel函数使用
    tiling.set_reduceAxisLen(reduceAxisLen);
    // 设置TilingData中除了reduceAxisLen之外其他成员变量的值
    ...
    // TilingData序列化保存
    tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
    context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
    ...
    return ge::GRAPH_SUCCESS;
}} // namespace optiling

编写核函数

核函数编程范式：

class KernelAdd {
public:
    __aicore__ inline KernelAdd(){}
    // 初始化函数，完成内存初始化相关操作
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){}
    // 核心处理函数，实现算子逻辑，调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作
    __aicore__ inline void Process(){}

private:
    // 搬入函数，完成CopyIn阶段的处理，被核心Process函数调用
    __aicore__ inline void CopyIn(int32_t progress){}
    // 计算函数，完成Compute阶段的处理，被核心Process函数调用
    __aicore__ inline void Compute(int32_t progress){}
    // 搬出函数，完成CopyOut阶段的处理，被核心Process函数调用
    __aicore__ inline void CopyOut(int32_t progress){}

private:
    TPipe pipe;  //Pipe内存管理对象
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;  //输入数据Queue队列管理对象，QuePosition为VECIN
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;  //输出数据Queue队列管理对象，QuePosition为VECOUT
    GlobalTensor<half> xGm, yGm, zGm;  //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出
};

编程过程中需要注意的有：

• 核函数必须具有void返回类型
• 使用__global__函数类型限定符来标识核函数
• 使用<<<…>>>调用核函数
• 使用__aicore__函数类型限定符来标识核函数在设备端AI Core上执行

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)

• 仅支持入参为指针类型或C/C++内置数据类型
• 使用__gm__类型限定符（指针变量指向global memory上的某处内存地址），可通过定义宏来避免过长类型名称
• 计算中用到的临时变量不能够入队出队，仅作临时计算使用
• 在重载算子实现时，优先调用开发者重载的算子

Device模块

Device模块负责指定计算运行的真实设备，简要介绍。

高级课程内容：矩阵编程与性能优化

矩阵乘法

矩阵编程学习过程中需要关注的概念：single、step、depth、iterator、ND/NZ。概念详细见
更多内容（同目录下还有融合算子编程）：AscendC开发文档-矩阵编程

性能优化

更多内容：AscendC开发文档-性能优化

总结

今天的课程中，从Ascend的研发需求和特点开始，深入到底层的算子计算过程系统学习了模型计算原理。实践方面，初步学习了Ascend C的编程方法，然后通过动手实践，初步了解了算子的编程方法。