本节以一个简单算子为例,带您体验从算子工程创建、代码编写、编译部署到运行验证的开发全流程,让您对算子开发工程有个宏观的认识,此处我们以输入是动态shape的Add算子实现为例,为了与内置Add算子区分,定义算子类型为AddCustom。
CANN软件包中提供了工程创建工具msOpGen,开发者可以输入算子原型定义文件生成Ascend C算子开发工程。
假设AddCustom算子的原型定义文件命名为add_custom.json,存储路径为: $HOME/sample,文件内容如下:
[
{
"op": "AddCustom",
"language": "cpp",
"input_desc": [
{
"name": "x",
"param_type": "required",
"format": [
"ND"
],
"type": [
"float16"
]
},
{
"name": "y",
"param_type": "required",
"format": [
"ND"
],
"type": [
"float16"
]
}
],
"output_desc": [
{
"name": "z",
"param_type": "required",
"format": [
"ND"
],
"type": [
"float16"
]
}
]
}
]
${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
AI处理器的型号<soc_version>请通过如下方式获取:
其中:
基于同系列的AI处理器型号创建的算子工程,其基础功能(基于该工程进行算子开发、编译和部署)通用。
AddCustom ├── build.sh // 编译入口脚本 ├── cmake // 算子工程编译所需脚本及公共编译文件存放目录 ├── CMakeLists.txt // 算子工程构建过程配置文件 ├── CMakePresets.json // 编译配置项 ├── framework // AI框架适配时,算子插件实现文件目录 ├── op_host // host侧实现文件 │ ├── add_custom_tiling.h // 算子tiling定义文件 │ ├── add_custom.cpp // 算子原型注册、shape推导、信息库、tiling实现等内容文件 │ ├── CMakeLists.txt ├── op_kernel // kernel侧实现文件 │ ├── CMakeLists.txt │ ├── add_custom.cpp // 算子核函数实现文件 ├── scripts // 自定义算子工程打包相关脚本所在目录
上述目录结构中的粗体文件为后续算子开发过程中需要修改的文件,其他文件无需修改。
在工程存储目录的“AddCustom/op_kernel/add_custom.cpp”文件中实现算子的核函数,完整的样例代码您可以在add_custom.cpp中查看,下面介绍关键实现代码。
算子核函数实现代码的内部调用关系示意图如下:
由此可见除了Init函数完成初始化外,Process中完成了对流水任务:“搬入、计算、搬出”的调用,开发者可以重点关注三个流水任务的实现。
1 2 3 4 5 6 7 8 9 10 11 |
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) { // 获取Host侧传入的Tiling参数 GET_TILING_DATA(tiling_data, tiling); // 初始化算子类 KernelAdd op; // 算子类的初始化函数,完成内存初始化相关工作 op.Init(x, y, z, tiling_data.totalLength, tiling_data.tileNum); // 完成算子实现的核心逻辑 op.Process(); } |
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 |
#include "kernel_operator.h" constexpr int32_t BUFFER_NUM = 2; class KernelAdd { public: __aicore__ inline KernelAdd() {} // 初始化函数,完成内存初始化相关操作 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum) { // 使用获取到的TilingData计算得到blockLength(每个核上总计算数据大小)、tileNum(每个核上分块个数)、tileLength(每个分块大小)等变量 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内存管理对象为输入输出Queue分配内存 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)); } // 核心处理函数,实现算子逻辑,调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作 __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: // 搬入函数,完成CopyIn阶段的处理,被核心Process函数调用 __aicore__ inline void CopyIn(int32_t progress) { // 从Queue中分配输入Tensor AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>(); AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>(); // 将GlobalTensor数据拷贝到LocalTensor AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength); AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength); // 将LocalTesor放入VECIN(代表矢量编程中搬入数据的逻辑存放位置)的Queue中 inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal); } // 计算函数,完成Compute阶段的处理,被核心Process函数调用 __aicore__ inline void Compute(int32_t progress) { // 将Tensor从队列中取出,用于后续计算 AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>(); AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>(); // 从Queue中分配输出Tensor AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>(); // 调用Add接口进行计算 AscendC::Add(zLocal, xLocal, yLocal, this->tileLength); // 将计算结果LocalTensor放入到VecOut的Queue中 outQueueZ.EnQue<DTYPE_Z>(zLocal); // 释放输入Tensor inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal); } // 搬出函数,完成CopyOut阶段的处理,被核心Process函数调用 __aicore__ inline void CopyOut(int32_t progress) { // 从VecOut的Queue中取出输出Tensor AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>(); // 将输出Tensor拷贝到GlobalTensor中 AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength); // 将不再使用的LocalTensor释放 outQueueZ.FreeTensor(zLocal); } private: //Pipe内存管理对象 AscendC::TPipe pipe; //输入数据Queue队列管理对象,TPosition为VECIN AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY; //输出数据Queue队列管理对象,TPosition为VECOUT AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ; //管理输入输出Global Memory内存地址的对象,其中xGm, yGm为输入,zGm为输出 AscendC::GlobalTensor<DTYPE_X> xGm; AscendC::GlobalTensor<DTYPE_Y> yGm; AscendC::GlobalTensor<DTYPE_Z> zGm; // 每个核上总计算数据大小 uint32_t blockLength; // 每个核上总计算数据分块个数 uint32_t tileNum; // 每个分块大小 uint32_t tileLength; }; |
核函数开发完成后,下一步就是进行Host侧的实现,对应“AddCustom/op_host”目录下的add_custom_tiling.h文件与add_custom.cpp文件。下面简要介绍下两个文件的关键实现,完整的样例代码可参见add_custom_tiling.h与add_custom.cpp。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 |
#ifndef ADD_CUSTOM_TILING_H #define ADD_CUSTOM_TILING_H #include "register/tilingdata_base.h" namespace optiling { BEGIN_TILING_DATA_DEF(TilingData) // AddCustom算子使用了2个tiling参数:totalLength与tileNum TILING_DATA_FIELD_DEF(uint32_t, totalLength); // 总计算数据量 TILING_DATA_FIELD_DEF(uint32_t, tileNum); // 每个核上总计算数据分块个数 END_TILING_DATA_DEF; // 注册tiling数据到对应的算子 REGISTER_TILING_DATA_CLASS(AddCustom, TilingData) } #endif // ADD_CUSTOM_TILING_H |
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 |
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 optiling |
Add算子的输出shape等于输入shape,所以直接将输入shape赋给输出shape,当前msOpGen工具生成的代码“InferShape”函数无需修改。
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 |
namespace ops { class AddCustom : public OpDef { public: explicit AddCustom(const char* name) : OpDef(name) { // Add算子的第一个输入 this->Input("x") .ParamType(REQUIRED) // 代表输入必选 .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32 }) // 输入支持的数据类型 .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND }); // 输入支持的数据格式 // Add算子的第二个输入 this->Input("y") .ParamType(REQUIRED) .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32 }) .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND }); // Add算子的输出 this->Output("z") .ParamType(REQUIRED) .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32 }) .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND }); // 关联InferShape函数 this->SetInferShape(ge::InferShape); // 关联Tiling函数 this->AICore() .SetTiling(optiling::TilingFunc); // 注册算子支持的AI处理器型号,请替换为实际支持的AI处理器型号 this->AICore().AddConfig("ascendxxx"); } }; // 结束算子注册 OP_ADD(AddCustom); } // namespace ops |
编译AddCustom工程,生成自定义算子安装包,并将其安装到算子库中。
./build.sh
编译成功后,会在当前目录下创建build_out目录,并在build_out目录下生成自定义算子安装包custom_opp_<target os>_<target architecture>.run,例如“custom_opp_ubuntu_x86_64.run”。
在自定义算子包所在路径下,执行如下命令,安装自定义算子包。
./custom_opp_<target os>_<target architecture>.run
命令执行成功后,自定义算子包中的相关文件将部署至当前环境的OPP算子库的vendors/customize目录中,如果用户部署多个自定义算子包,可通过如下命令指定路径安装:
./custom_opp_<target os>_<target architecture>.run --install-path=<path>
说明:如果部署算子包时通过配置--install-path参数指定了算子包的安装目录,则在使用自定义算子前,需要执行source <path>/vendors/<vendor_name>/bin/set_env.bash命令,set_env.bash脚本中将自定义算子包的安装路径追加到环境变量ASCEND_CUSTOM_OPP_PATH中,使自定义算子在当前环境中生效。
├── opp // 算子库目录
│ ├── built-in // 内置算子所在目录
│ ├── vendors // 自定义算子所在目录
│ ├── config.ini
│ └── vendor_name1 // 自定义算子所在目录,若不指定路径安装,默认为“customize”
│ ├── framework //自定义算子插件库
│ ├── op_impl
│ │ └── ai_core
│ │ └── tbe
│ │ ├── config
│ │ │ └── ${soc_version} //昇腾AI处理器类型
│ │ │ └── aic-${soc_version}-ops-info.json //自定义算子信息库文件
│ │ ├── vendor_name1_impl //自定义算子实现代码文件
│ │ │ └── dynamic
│ │ │ ├── xx.cpp
│ │ │ └── xx.py
│ │ ├── kernel //自定义算子二进制文件
│ │ │ └── ${soc_version} //昇腾AI处理器类型
│ │ │ └── config
│ │ └── op_tiling
│ │ ├── lib
│ │ └── liboptiling.so
│ └── op_proto //自定义算子原型库所在目录
│ ├── inc
│ │ └── op_proto.h
│ └── lib
│ ├── vendor_name2 // 存储厂商vendor_name2部署的自定义算子
CANN开发套件包中提供了ST测试工具“msOpST”,用于生成算子的ST测试用例并在硬件环境中执行。
本节仅以AddCustom算子为例,介绍ST测试工具的关键执行流程。
[
{
"case_name": "Test_AddCustom_001",
"op": "AddCustom",
"input_desc": [
{
"format": [
"ND"
],
"type": [
"float16"
],
"shape": [8,2048],
"data_distribute": [
"uniform"
],
"value_range": [
[
0.1,
1.0
]
],
"name": "x"
},
{
"format": [
"ND"
],
"type": [
"float16"
],
"shape": [8,2048],
"data_distribute": [
"uniform"
],
"value_range": [
[
0.1,
1.0
]
],
"name": "y"
}
],
"output_desc": [
{
"format": [
"ND"
],
"type": [
"float16"
],
"shape": [8,2048],
"name": "z"
}
]
}
]
export DDK_PATH=${INSTALL_DIR} export NPU_HOST_LIB=${INSTALL_DIR}/{arch-os}/devlib
提示:请根据CANN软件包实际安装路径对以上环境变量进行修改。
cd $HOME/Ascend/ascend-toolkit/latest/python/site-packages/bin
./msopst run -i $HOME/AddCustom_st/AddCustom_case.json -soc <soc_version> -out $HOME/AddCustom_st
此命令执行完成后,会输出类似如下打屏结果:
1 2 3 4 5 6 7 |
------------------------------------------------------------------------ - test case count: 1 - success count: 1 - failed count: 0 ------------------------------------------------------------------------ 2023-08-28 20:20:40 (25058) - [INFO] Process finished! 2023-08-28 20:20:40 (25058) - [INFO] The st report saved in: xxxx/AddCustom_st/20230828202015/st_report.json. |
您也可以查看上述屏显信息提示的“st_report.json”文件,查看详细运行结果。