当前位置：首页 > news >正文

OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用 - Tiling实现 2

news 来源：原创 2024/9/19 9:37:37

OrangePi AIpro 香橙派昇腾 Ascend C 算子开发与调用 - Tiling实现 2

flyfish

前置知识 1
前置知识 2

Host侧CPU和Device侧NPU的主要区别

不同的硬件资源

CPU是为了执行通用计算任务而设计的，但在处理大量的并行计算（如矩阵乘、批数据处理）时效率不高。NPU是为了加速机器学习和深度学习任务而设计的，它擅长执行大量的并行计算。NPU包含了大量的专用硬件，比如：支持矩阵计算的Cube单元，NPU中一个核可以支持一个时钟周期内完成数据量为161616、数据类型为fp16的矩阵乘法；支持向量计算的Vector单元，NPU中一个核可以支持一个时钟周期内处理128个fp16的加法。

不同的物理内存空间

Host和Device的物理内存是分离的，有时需要在Host侧内存和Device侧内存之间进行数据交换。

Ascend C的算子实现主要包含两个部分：

Host侧Tiling实现

由于NPU中AI Core内部存储无法完全容纳算子输入输出的所有数据，需要每次搬运一部分输入数据进行计算然后搬出，再搬运下一部分输入数据进行计算，这个过程就称之为Tiling。切分数据的算法称为Tiling算法或者Tiling策略。根据算子的shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循环多少次）的计算程序，称之为Tiling实现，也叫Tiling函数（Tiling Function）。由于Tiling实现中完成的均为标量计算，AI Core并不擅长，所以我们将其独立出来放在Host侧CPU上执行。

Device侧Kernel实现

Kernel实现即算子核函数实现，在Kernel函数内部通过解析Host侧传入的Tiling结构体获取Tiling信息，根据Tiling信息控制数据搬入搬出Local Memory的流程；通过调用计算、数据搬运、内存管理、任务同步API，实现算子逻辑。其核心逻辑基本上都为计算密集型任务，需要在NPU上执行。

算子Tiling传递数据流

算子数据流
算子执行过程中涉及到Host和Device的数据交换。这里仅针对Tiling参数的传递，给出具体的数据流：Host侧Tiling算法根据算子具体输入输出的信息，完成Tiling参数的计算，存放在Tiling结构体中；将Host侧的Tiling结构体发送到Device侧，Device侧的算子获取并解析Tiling结构体，基于该信息执行后续的算子计算逻辑。
**加粗样式**

完整的代码放置最后
AddCustom算子的原型定义文件命名为add_custom.json

[{"op": "AddCustom","input_desc": [{"name": "x","param_type": "required","format": ["ND"],"type": ["fp16"]},{"name": "y","param_type": "required","format": ["ND"],"type": ["fp16"]}],"output_desc": [{"name": "z","param_type": "required","format": ["ND"],"type": ["fp16"]}]}
]

${INSTALL_DIR}/python/site-packages/bin/msopgen gen -i $HOME/sample/add_custom.json -c ai_core-<soc_version> -lan cpp -out   $HOME/sample/AddCustom

${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际环境进行替换。
-i：算子原型定义文件add_custom.json所在路径。
-c：ai_core-<soc_version>代表算子在AI Core上执行，<soc_version>为昇腾AI处理器的型号，可通过npu-smi info命令进行查询，基于同系列的AI处理器型号创建的算子工程，其基础功能能力通用。
例如soc_version设置为Ascend310P1，创建的算子工程，也可以用于开发运行于Ascend310P3上的算子。
-lan： 参数cpp代表算子基于Ascend C编程框架，使用C++编程语言开发。

实际执行的命令，截断显示多行是为了看清楚
add_custom.json所在路径是/home/HwHiAiUser/sample/add_custom.json

/usr/local/Ascend/ascend-toolkit/8.0.RC3.alpha002/python/site-packages/bin/msopgen gen 
-i /home/HwHiAiUser/sample/add_custom.json
-c ai_core-Ascend310B -lan cpp 
-out   /home/HwHiAiUser/sample/AddCustom

执行命令后生成的文件
请添加图片描述

在这里插入图片描述

AddCustom
├── build.sh         // 编译入口脚本
├── cmake 
│   ├── config.cmake
│   ├── func.cmake
│   ├── intf.cmake
│   ├── makeself.cmake
│   └── util        // 算子工程编译所需脚本及公共编译文件存放目录
├── CMakeLists.txt    // 算子工程的CMakeLists.txt
├── CMakePresets.json    // 编译配置项
├── framework        // 算子插件实现文件目录，单算子模型文件的生成不依赖算子适配插件，无需关注
├── op_host                      // host侧实现文件
│   ├── add_custom_tiling.h    // 算子tiling定义文件
│   ├── add_custom.cpp         // 算子原型注册、shape推导、信息库、tiling实现等内容文件
│   ├── CMakeLists.txt
├── op_kernel                   // kernel侧实现文件
│   ├── CMakeLists.txt   
│   ├── add_custom.cpp        // 算子代码实现文件 
└── scripts                     // 自定义算子工程打包相关脚本所在目录

替换为CANN软件包安装后的实际路径
修改CMakePresets.json中ASCEND_CANN_PACKAGE_PATH为CANN软件包安装路径在这里插入图片描述
如果不改直接编译./build.sh就出现如下问题

在这里插入图片描述
在Ascend C中，Tiling策略的直接表示形式是一个C语言中的结构体(struct)，简称Tiling结构体
Tiling结构体定义在Tiling头文件(形如xxxx custom tiling.h)中，其中的每个结构体参数表示如何对输入数据进行切分，以及决定了计算过程的一些细节，结构体在host侧实例化，并通过指针传入kernel函数中

AddCustom/op_host/add_custom_tiling.h

#include "register/tilingdata_base.h"namespace optiling {
BEGIN_TILING_DATA_DEF(AddCustomTilingData)TILING_DATA_FIELD_DEF(uint32_t, size);
END_TILING_DATA_DEF;REGISTER_TILING_DATA_CLASS(AddCustom, AddCustomTilingData)
}

请添加图片描述代码框架编写：
需要增加#ifndef…的判断条件，防止头文件的重复包含；需要包含register/tilingdata_base.h头文件，tilingdata_base.h中定义了多个用于tilingdata注册的宏。

TilingData参数设计，TilingData参数本质上是和并行数据切分相关的参数，本示例算子使用了2个tiling参数：totalLength、tileNum。totalLength是指需要计算的数据量大小，tileNum是指每个核上总计算数据分块个数。比如，totalLength这个参数传递到kernel侧后，可以通过除以参与计算的核数，得到每个核上的计算量，这样就完成了多核数据的切分。
TilingData结构定义，通过BEGIN_TILING_DATA_DEF接口定义一个TilingData的类，通过TILING_DATA_FIELD_DEF接口增加TilingData的两个字段totalLength、tileNum，通过END_TILING_DATA_DEF接口结束TilingData定义。

BEGIN_TILING_DATA_DEF(TilingData)               // 注册一个tiling的类，以tiling的名字作为入参TILING_DATA_FIELD_DEF(uint32_t, totalLength); // 添加tiling字段，总计算数据量TILING_DATA_FIELD_DEF(uint32_t, tileNum);     // 添加tiling字段，每个核上总计算数据分块个数
END_TILING_DATA_DEF;

注册TilingData结构，通过REGISTER_TILING_DATA_CLASS接口，注册TilingData类，和自定义算子相关联。REGISTER_TILING_DATA_CLASS第一个参数为op_type(算子类型)，本样例中传入AddCustom，第二个参数为TilingData的类名。

注册算子tilingdata类到对应的AddCustom算子REGISTER_TILING_DATA_CLASS(AddCustom, TilingData)

AddCustom/op_host/add_custom.cpp

#include "add_custom_tiling.h"
#include "register/op_def_registry.h"namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{AddCustomTilingData tiling;const gert::StorageShape* x1_shape = context->GetInputShape(0);int32_t data_sz = 1;for (int i = 0; i < x1_shape->GetStorageShape().GetDimNum(); i++)data_sz *= x1_shape->GetStorageShape().GetDim(i);tiling.set_size(data_sz);context->SetBlockDim(8);tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());return ge::GRAPH_SUCCESS;
}
}

在这里插入图片描述

请添加图片描述

AddCustom/op_kernel/add_custom.cpp

Kernel侧使用Tiling信息
Kernel侧需要接收Tiling信息时，核函数定义是这样的:

__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling))

注意这里参数的顺序按照“输入、输出、workspace、tiling”的顺序排布，不要不要调整其顺序

#include "kernel_operator.h"
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {GET_TILING_DATA(tiling_data, tiling);// TODO: user kernel impl
}

请添加图片描述

固定shape与动态shape 区别

文件名	职能	固定shape	动态shape
main.cpp	主机侧数据文件读写以及主机侧到设备侧的数据拷贝，任务下发以及同步等待等	读取输入参数和申请内存，调用核函数等	新增tiling参数的内存申请，搬运与释放逻辑
add_custom.cpp	Ascend C算子核函数的实现	shape等参数以常量展现，编译期已知	shape等参数以入参展现，编译期未知
add_custom.py	输入数据和真值数据的生成	生成输入数据x和y，真值数据golden	新增生成tiling的.bin数据文件
CMakeLists.txt	管理工程项目编译构建配置	无变化	无变化
data_utils.h	主机侧数据打印等辅助函数的实现	无变化	无变化
run.sh	集成算子运行一体化脚本	无变化	无变化
add_custom_tiling.h	定义动态shape的tiling配置	不涉及	Tiling结构体与解析tiling宏函数

AddCustom/op_host/add_custom_tiling.h

#ifndef ADD_CUSTOM_TILING_H
#define ADD_CUSTOM_TILING_H
#include "register/tilingdata_base.h"namespace optiling {
BEGIN_TILING_DATA_DEF(TilingData)
TILING_DATA_FIELD_DEF(uint32_t, totalLength);
TILING_DATA_FIELD_DEF(uint32_t, tileNum);
END_TILING_DATA_DEF;REGISTER_TILING_DATA_CLASS(AddCustom, TilingData)
} // namespace optiling
#endif // ADD_CUSTOM_TILING_H

AddCustom/op_host/add_custom.cpp

#include "add_custom_tiling.h"
#include "register/op_def_registry.h"namespace optiling {
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{AddCustomTilingData tiling;const gert::StorageShape* x1_shape = context->GetInputShape(0);int32_t data_sz = 1;for (int i = 0; i < x1_shape->GetStorageShape().GetDimNum(); i++)data_sz *= x1_shape->GetStorageShape().GetDim(i);tiling.set_size(data_sz);context->SetBlockDim(8);tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());return ge::GRAPH_SUCCESS;
}
}namespace ge {
static ge::graphStatus InferShape(gert::InferShapeContext* context)
{const gert::Shape* x1_shape = context->GetInputShape(0);gert::Shape* y_shape = context->GetOutputShape(0);*y_shape = *x1_shape;return GRAPH_SUCCESS;
}
}namespace ops {
class AddCustom : public OpDef {
public:explicit AddCustom(const char* name) : OpDef(name){this->Input("x").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND});this->Input("y").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND});this->Output("z").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND});this->SetInferShape(ge::InferShape);this->AICore().SetTiling(optiling::TilingFunc);this->AICore().AddConfig("ascend310b");}
};OP_ADD(AddCustom);
}

更多配置的版本

#include "add_custom_tiling.h"
#include "register/op_def_registry.h"namespace optiling {
const uint32_t BLOCK_DIM = 8;
const uint32_t TILE_NUM = 8;
static ge::graphStatus TilingFunc(gert::TilingContext *context)
{TilingData tiling;uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();context->SetBlockDim(BLOCK_DIM);tiling.set_totalLength(totalLength);tiling.set_tileNum(TILE_NUM);tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());size_t *currentWorkspace = context->GetWorkspaceSizes(1);currentWorkspace[0] = 0;return ge::GRAPH_SUCCESS;
}
} // namespace optilingnamespace ge {
static graphStatus InferShape(gert::InferShapeContext *context)
{const gert::Shape *x1_shape = context->GetInputShape(0);gert::Shape *y_shape = context->GetOutputShape(0);*y_shape = *x1_shape;return GRAPH_SUCCESS;
}static graphStatus InferDataType(gert::InferDataTypeContext *context)
{const auto inputDataType = context->GetInputDataType(0);context->SetOutputDataType(0, inputDataType);return ge::GRAPH_SUCCESS;
}
} // namespace genamespace ops {
class AddCustom : public OpDef {
public:explicit AddCustom(const char *name) : OpDef(name){this->Input("x").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND});this->Input("y").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND});this->Output("z").ParamType(REQUIRED).DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND});this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);this->AICore().SetTiling(optiling::TilingFunc).AddConfig("ascend910").AddConfig("ascend310p").AddConfig("ascend310b").AddConfig("ascend910b");}
};
OP_ADD(AddCustom);
} // namespace ops

AddCustom/op_kernel/add_custom.cpp

#include "kernel_operator.h"
constexpr int32_t BUFFER_NUM = 2; // tensor num for each queueclass KernelAdd {
public:__aicore__ inline KernelAdd() {}__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum){this->blockLength = totalLength / AscendC::GetBlockNum();this->tileNum = tileNum;this->tileLength = this->blockLength / tileNum / BUFFER_NUM;xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);zGm.SetGlobalBuffer((__gm__ DTYPE_Z *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z));}__aicore__ inline void Process(){int32_t loopCount = this->tileNum * BUFFER_NUM;for (int32_t i = 0; i < loopCount; i++) {CopyIn(i);Compute(i);CopyOut(i);}}private:__aicore__ inline void CopyIn(int32_t progress){AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>();AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);inQueueX.EnQue(xLocal);inQueueY.EnQue(yLocal);}__aicore__ inline void Compute(int32_t progress){AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>();AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>();AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);outQueueZ.EnQue<DTYPE_Z>(zLocal);inQueueX.FreeTensor(xLocal);inQueueY.FreeTensor(yLocal);}__aicore__ inline void CopyOut(int32_t progress){AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);outQueueZ.FreeTensor(zLocal);}private:AscendC::TPipe pipe;AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;AscendC::GlobalTensor<DTYPE_X> xGm;AscendC::GlobalTensor<DTYPE_Y> yGm;AscendC::GlobalTensor<DTYPE_Z> zGm;uint32_t blockLength;uint32_t tileNum;uint32_t tileLength;
};extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{GET_TILING_DATA(tiling_data, tiling);KernelAdd op;op.Init(x, y, z, tiling_data.totalLength, tiling_data.tileNum);op.Process();
}#ifndef ASCENDC_CPU_DEBUG
// call of kernel function
void add_custom_do(uint32_t blockDim, void *l2ctrl, void *stream, uint8_t *x, uint8_t *y, uint8_t *z,uint8_t *workspace, uint8_t *tiling)
{add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z, workspace, tiling);
}
#endif

Tiling函数是在host侧实现的代码，与之相对应的，在kernel侧算子实现代码中，只需通过调用GET_TILING_DATA即可获取TilingData结构体参数，并使用具体的参数进行后续的计算。

通过Add算子举例来说明，固定shape和动态shape kernel侧算子实现的区别。

固定shape的算子例子中，TILE_NUM（每个核上总计算数据分块个数）、BLOCK_LENGTH（每个核上总计算数据大小）、TILE_LENGTH（每个分块大小）等是固定的数值。

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; // each tile length is seperated to 2 part, due to double buffer

动态shape的实现中，需要在核函数中通过GET_TILING_DATA获取Tiling参数，再基于Tiling参数计算得到singleCoreSize（每个核上总计算数据大小）、tileNum（每个核上总计算数据分块个数）、singleTileLength（每个分块大小）等变量。注意，对应的算子host实现中需要定义TilingData结构体，实现并注册计算TilingData的Tiling函数。具体请参考Tiling实现。
核函数中调用GET_TILING_DATA获取TilingData的样例如下：

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{GET_TILING_DATA(tilingData, tiling);KernelAdd op;op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);if (TILING_KEY_IS(1)) {op.Process();}
}

算子类的Init函数中，使用获取到的TilingData计算得到singleCoreSize、tileNum、singleTileLength等变量的样例如下。

__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
{ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");this->blockLength = totalLength / GetBlockNum();this->tileNum = tileNum;ASSERT(tileNum != 0 && "tile num can not be zero!");this->tileLength = this->blockLength / tileNum / BUFFER_NUM;// ...}