北研-鲲鹏&昇腾创新大赛-Ascend C算子开发笔记

AI Core内部并行计算架构抽象

AI Core中包含计算单元、存储单元、搬运单元等核心组件

• 计算单元：Scalar计算单元、Cube计算单元、Vector计算单元
• 存储单元：AI Core的内部存储统称为Local Memory，外部存储为Global Memory
• 搬运单元：负责在Global Memory和Local Memory之间搬运数据，包含搬运单元MTE2（数据搬入单元），MTE3（数据搬出单元）

算子基本概念

• 算子名称（Name）：用于标志网络中的某一个算子，同一网络中的算子名称需要保持唯一e.g.：卷积算子Conv1，Conv2
• 算子类型（Type）：同类型的算子的实现逻辑相同
• 数据容器（Tensor）：张量（Tensor）是存储算子输入数据、输出数据的容器，而张量描述符（TensoeDesc）是对输入数据、输出数据的描述。数据结构如下

• 数据排布格式（format）：

  • 深度学习领域，多维数据通过多维数组存储，比如卷积神经网络用四维数组
    • N：Batch数量，比如图像的数目
    • H：特征图高度
    • W：特征图宽度
    • C：channel
  • Caffe->NCHW
• 数据类型（dtype）：FP16...
• 形状（Shape）：比如（10，）、（1024，1024）、（2，3，4）、（i1，i2，i3...）
• 名称（Name）：用于索引，保持唯一
• 算子属性
  1. 轴——表示数据维度

什么是核函数

核函数是Ascend C算子设备侧的入口

Ascend C    __global__ __aicore__ void kernel_name(argument list);

CUDA         __global__ void kernel_name(argument list);

核函数是直接在设备侧执行的代码。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，SPMD编程模型允许核函数调用时，多个核并行的执行同一个计算任务。

如何编写核函数

使用函数类型限定符

除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__global__和__aicore__

使用__global__函数类型限定符来标识他是一个核函数，可以被<<<...>>>调用；使用__aicore__函数类型限定符来标识该和函数在设备侧AI Core上执行：

__global__ __aicore__ void kernel_name(argument list);

SPMD编程模型

Ascend C算子编程是SPMD的编程，将需要处理的数据拆分并分布在多个计算核心上运行多个AI Core共享相同的指令代码，每个核上的运行实例唯一的区别是block_idx不同block的类似于进程，block_idx就是标识进程唯一性的进程ID，编程中使用函数GetBlockIdx()获取ID。

SPMD数据并行示意图

SPMD并行计算示意图

sinh算子实现

Ascend C sinh算子设计规格

算子类型（OpType）	SinhCustom
算子输入	name	shape	data type	format
	x	(8, 2048)	half	ND
算子输出	output_z	(8, 2048)	half	ND
核函数名称	sinh_custom
使用的主要接口	DataCopy：数据搬移接口
	Exp：单目指数运算
	Muls：矢量标量相乘
	Sub：矢量相减
算子实现文件名称	sinh_custom.cpp

 Ascend C中kernel侧中compute实现

        LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();
        //Add(zLocal, xLocal, yLocal, this->tileLength);
        Exp(yLocal,xLocal,this->tileLength);
        Muls(xLocal,xLocal,(half)(-1),this->tileLength);
        Exp(xLocal,xLocal,this->tileLength);
        Sub(yLocal,yLocal,xLocal,this->tileLength);
        Muls(yLocal,yLocal,(half)0.5,this->tileLength);
        outQueueY.EnQue<DTYPE_Y>(yLocal);
        inQueueX.FreeTensor(xLocal);

 host侧tiling结构体定义和注册

#ifndef SINH_CUSTOM_TILING_H
#define SINH_CUSTOM_TILING_H
#include "register/tilingdata_base.h"

namespace optiling {
BEGIN_TILING_DATA_DEF(SinhCustomTilingData)
  TILING_DATA_FIELD_DEF(uint32_t, totalLength);
  TILING_DATA_FIELD_DEF(uint32_t, tileNum);
END_TILING_DATA_DEF;

REGISTER_TILING_DATA_CLASS(SinhCustom, SinhCustomTilingData)
}
#endif

 先编写Kernel，可以先进行CPU侧调试，在进行NPU侧调试。

bash run.sh -r cpu -v ascend310B1   #cpu测试
bash run.sh -r npu -v ascend310B1   #npu测试

算子工程编译

SinhCustom/SinhCustom/build.sh
SinhCustom/SinhCustom/build_out/custom_opp_euleros_aarch64.run

aclnn验证

bash SinhCustom/AclNNInvocation/run.sh

 通过后显示即通过sinh用例