时间: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的编程方法,然后通过动手实践,初步了解了算子的编程方法。