AscendC从入门到醒目系列（三）基于自界说算子工程开发AscendC算子 ...

星球的眼睛  金牌会员 | 2024-11-12 20:12:10 | 显示全部楼层 | 阅读模式
打印 上一主题 下一主题
楼主

主题 502|帖子 502|积分 1506

本次主要讨论下AscendC别的一种开发流程，基于自界说算子工程的算子开发。从算子工程创建、代码编写、编译部署到运行验证的开发全流程，让您对算子开发工程有个宏观的熟悉，此处我们以输入是动态shape（主要体如今tiling）的Add算子实现为例，为了与内置Add算子区分，界说算子范例为AddCustom。
1、创建工程

CANN软件包中提供了工程创建工具msOpGen，开发者可以输入算子原型界说文件天生Ascend C算子开发工程。
1.1 编写AddCustom算子的原型界说json文件

  2. ```yaml
  4. ```java
  5. [
  6.     {
  7.         "op": "AddCustom",
  8.         "input_desc": [
  9.             {
  10.                 "name": "x",
  11.                 "param_type": "required",
  12.                 "format": [
  13.                     "ND"
  14.                 ],
  15.                 "type": [
  16.                     "fp16"
  17.                 ]
  18.             },
  19.             {
  20.                 "name": "y",
  21.                 "param_type": "required",
  22.                 "format": [
  23.                     "ND"
  24.                 ],
  25.                 "type": [
  26.                     "fp16"
  27.                 ]
  28.             }
  29.         ],
  30.         "output_desc": [
  31.             {
  32.                 "name": "z",
  33.                 "param_type": "required",
  34.                 "format": [
  35.                     "ND"
  36.                 ],
  37.                 "type": [
  38.                     "fp16"
  39.                 ]
  40.             }
  41.         ]
  42.     }
  43. ]
复制代码
1.2 用msOpGen工具天生AddCustom算子的开发工程

  1. ${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软件安装后文件存储路径，请根据现真相况举行替换，如/usr/local/Ascend/ascend-toolkit/latest。
  • -i：算子原型界说文件add_custom.json所在路径。
  • -c：ai_core-<soc_version>代表算子在AI Core上执行，<soc_version>为昇腾AI处理器的型号，可通过npu-smi info命令举行查询，基于同系列的AI处理器型号创建的算子工程，其根本能力通用。例如soc_version设置为Ascend310P1，Ascend910B3等。
  • -lan： 参数cpp代表算子基于Ascend C编程框架，利用C++编程语言开发。
1.3 工程目录天生

命令执行完后，会在$HOME/sample目录下天生算子工程目录AddCustom，工程中包含算子实现的模板文件，编译脚本等，如下所示
  1. AddCustom
  2. ├── build.sh         // 编译入口脚本
  3. ├── cmake
  4. │   ├── config.cmake
  5. │   ├── util        // 算子工程编译所需脚本及公共编译文件存放目录
  6. ├── CMakeLists.txt   // 算子工程的CMakeLists.txt
  7. ├── CMakePresets.json // 编译配置项
  8. ├── framework        // 算子插件实现文件目录，单算子模型文件的生成不依赖算子适配插件，无需关注
  9. ├── op_host                      // host侧实现文件
  10. │   ├── add_custom_tiling.h    // 算子tiling定义文件
  11. │   ├── add_custom.cpp         // 算子原型注册、shape推导、信息库、tiling实现等内容文件
  12. │   ├── CMakeLists.txt
  13. ├── op_kernel                   // kernel侧实现文件
  14. │   ├── CMakeLists.txt   
  15. │   ├── add_custom.cpp        // 算子核函数实现文件
  16. ├── scripts                     // 自定义算子工程打包相关脚本所在目录
复制代码


  • CMakePresets.json // 编译设置项
  • add_custom_tiling.h // 算子tiling界说文件
  • op_host/add_custom.cpp // 算子原型注册、shape推导、信息库、tiling实现等内容文件
  • op_kernel/add_custom.cpp // 算子核函数实现文件
    上述文件为后续算子开发过程中需要修改的文件，其他文件无需修改。
2 算子核函数实现

在工程存储目录的“AddCustom/op_kernel/add_custom.cpp”文件中实现算子的核函数。算子核函数实当代码的内部调用关系示意图如下：

2.1 核函数界说

核函数的界说，并在核函数中调用算子类的Init和Process函数。
  1. extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
  2. {
  3.     // 获取Host侧传入的Tiling参数
  4.     GET_TILING_DATA(tiling_data, tiling);
  5.     // 初始化算子类
  6.     KernelAdd op;
  7.     // 算子类的初始化函数，完成内存初始化相关工作
  8.     op.Init(x, y, z, tiling_data.totalLength, tiling_data.tileNum);
  9.     // 完成算子实现的核心逻辑
  10.     op.Process();
  11. }
复制代码
2.2 界说KernelAdd算子类

和之前AscendC从入门到醒目系列（二） - 知乎 (zhihu.com)中一样，KernelAdd算子类主要也是实现Init，CopyIn，Compute，CopyOut这个4个关键函数。
  1. #include "kernel_operator.h"
  2. constexpr int32_t BUFFER_NUM = 2;
  3. class KernelAdd {
  4. public:
  5.     __aicore__ inline KernelAdd() {}
  6.     // 初始化函数，完成内存初始化相关操作
  7.     __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
  8.     {
  9.         // 使用获取到的TilingData计算得到singleCoreSize(每个核上总计算数据大小)、tileNum（每个核上分块个数）、singleTileLength（每个分块大小）等变量
  10.         this->blockLength = totalLength / AscendC::GetBlockNum();
  11.         this->tileNum = tileNum;
  12.         this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
  13.         
  14.         // 获取当前核的起始索引
  15.         xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
  16.         yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
  17.         zGm.SetGlobalBuffer((__gm__ DTYPE_Z*)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
  18.         // 通过Pipe内存管理对象为输入输出Queue分配内存
  19.         pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
  20.         pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));
  21.         pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Z));
  22.     }
  23.     // 核心处理函数，实现算子逻辑，调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作
  24.     __aicore__ inline void Process()
  25.     {
  26.         int32_t loopCount = this->tileNum * BUFFER_NUM;
  27.         for (int32_t i = 0; i < loopCount; i++) {
  28.             CopyIn(i);
  29.             Compute(i);
  30.             CopyOut(i);
  31.         }
  32.     }
  35. private:
  36.     // 搬入函数，完成CopyIn阶段的处理，被核心Process函数调用
  37.     __aicore__ inline void CopyIn(int32_t progress)
  38.     {
  39.         // 从Queue中分配输入Tensor
  40.         AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
  41.         AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>();
  42.          // 将GlobalTensor数据拷贝到LocalTensor
  43.         AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
  44.         AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
  45.         // 将LocalTesor放入VECIN（代表矢量编程中搬入数据的逻辑存放位置）的Queue中
  46.         inQueueX.EnQue(xLocal);
  47.         inQueueY.EnQue(yLocal);
  48.     }
  49.     // 计算函数，完成Compute阶段的处理，被核心Process函数调用
  50.     __aicore__ inline void Compute(int32_t progress)
  51.     {
  52.         // 将Tensor从队列中取出，用于后续计算
  53.         AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
  54.         AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>();
  55.         // 从Queue中分配输出Tensor
  56.         AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>();
  57.         // 调用Add接口进行计算
  58.         AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);
  59.         // 将计算结果LocalTensor放入到VecOut的Queue中
  60.         outQueueZ.EnQue<DTYPE_Z>(zLocal);
  61.         // 释放输入Tensor
  62.         inQueueX.FreeTensor(xLocal);
  63.         inQueueY.FreeTensor(yLocal);
  64.     }
  65.     // 搬出函数，完成CopyOut阶段的处理，被核心Process函数调用
  66.     __aicore__ inline void CopyOut(int32_t progress)
  67.     {
  68.         // 从VecOut的Queue中取出输出Tensor
  69.         AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();
  70.         // 将输出Tensor拷贝到GlobalTensor中
  71.         AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
  72.         // 将不再使用的LocalTensor释放
  73.         outQueueZ.FreeTensor(zLocal);
  74.     }
  77. private:
  78.     //Pipe内存管理对象
  79.     AscendC::TPipe pipe;
  80.     //输入数据Queue队列管理对象，QuePosition为VECIN
  81.     AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
  82.     //输出数据Queue队列管理对象，QuePosition为VECOUT
  83.     AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
  84.     //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出
  85.     AscendC::GlobalTensor<DTYPE_X> xGm;
  86.     AscendC::GlobalTensor<DTYPE_Y> yGm;
  87.     AscendC::GlobalTensor<DTYPE_Z> zGm;
  88.     // 每个核上总计算数据大小
  89.     uint32_t blockLength;
  90.     // 每个核上总计算数据分块个数
  91.     uint32_t tileNum;
  92.     // 每个分块大小
  93.     uint32_t tileLength;
  94. };
复制代码
3. host侧开发

核函数开发并验证完成后，下一步就是举行Host侧的实现，对应“AddCustom/op_host”目录下的add_custom_tiling.h文件与add_custom.cpp文件。
3.1 add_custom_tiling.h

这个是界说数据怎么切分，每个核上执行多少数据量，核上的数据又怎么切分执行的问题。
  1. #ifndef ADD_CUSTOM_TILING_H
  2. #define ADD_CUSTOM_TILING_H
  3. #include "register/tilingdata_base.h"
  4. namespace optiling {
  5. BEGIN_TILING_DATA_DEF(TilingData)
  6.   // AddCustom算子使用了2个tiling参数：totalLength与tileNum
  7.   TILING_DATA_FIELD_DEF(uint32_t, totalLength);     // 总计算数据量
  8.   TILING_DATA_FIELD_DEF(uint32_t, tileNum);         // 每个核上总计算数据分块个数
  9. END_TILING_DATA_DEF;
  11. // 注册tiling数据到对应的算子
  12. REGISTER_TILING_DATA_CLASS(AddCustom, TilingData)
  13. }
  14. #endif // ADD_CUSTOM_TILING_H
复制代码
3.2 add_custom.cpp

修改“add_custom.cpp”文件，举行Tiling的实现。
  1. namespace optiling {
  2. const uint32_t BLOCK_DIM = 8;
  3. const uint32_t TILE_NUM = 8;
  4. static ge::graphStatus TilingFunc(gert::TilingContext* context)
  5. {
  6.     TilingData tiling;
  7.     uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();
  8.     context->SetBlockDim(BLOCK_DIM);
  9.     tiling.set_totalLength(totalLength);
  10.     tiling.set_tileNum(TILE_NUM);
  11.     tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
  12.     context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
  13.     size_t *currentWorkspace = context->GetWorkspaceSizes(1);
  14.     currentWorkspace[0] = 0;
  15.     return ge::GRAPH_SUCCESS;
  16. }
  17. } // namespace optiling
复制代码
3.3 实现AddCustom算子的shape推导

在“add_custom.cpp”文件中实现AddCustom算子的shape推导。
  1. static graphStatus InferShape(gert::InferShapeContext *context)
  2. {
  3.     const gert::Shape *x1_shape = context->GetInputShape(0);
  4.     gert::Shape *y_shape = context->GetOutputShape(0);
  5.     *y_shape = *x1_shape;
  6.     return GRAPH_SUCCESS;
  7. }
复制代码
3.4 算子原型注册

  1. namespace ops {
  2. class AddCustom : public OpDef {
  3. public:
  4.     explicit AddCustom(const char* name) : OpDef(name)
  5.     {
  6.         // Add算子的第一个输入
  7.         this->Input("x")
  8.             .ParamType(REQUIRED)    // 代表输入必选
  9.             .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32 })   // 输入支持的数据类型
  10.             .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });   // 输入支持的数据格式
  11.         // Add算子的第二个输入
  12.         this->Input("y")
  13.             .ParamType(REQUIRED)
  14.             .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32 })
  15.             .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });
  16.         this->Output("z")
  17.             .ParamType(REQUIRED)
  18.             .DataType({ ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_INT32 })
  19.             .Format({ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND });
  20.         // 关联InferShape函数
  21.         this->SetInferShape(ge::InferShape);
  22.         // 关联Tiling函数
  23.         this->AICore()
  24.             .SetTiling(optiling::TilingFunc);
  25.         // 注册算子支持的AI处理器型号，请替换为实际支持的AI处理器型号
  26.         this->AICore().AddConfig("ascendxxx");
  27.     }
  28. };
  29. // 结束算子注册
  30. OP_ADD(AddCustom);
  31. } // namespace ops
复制代码
OP_ADD(AddCustom)：算子原型注册接口 .
4 算子工程编译部署

编译AddCustom工程，天生自界说算子安装包，并将其安装到算子库中
4.1 编译自界说算子工程

编译自界说算子工程，构建天生自界说算子包。修改CMakePresets.json中ASCEND_CANN_PACKAGE_PATH为CANN软件的安装目录，例如：/usr/local/Ascend/ascend-toolkit/latest。
  1. {
  2.     ……
  3.     "configurePresets": [
  4.         {
  5.                 ……
  6.                 "ASCEND_CANN_PACKAGE_PATH": {
  7.                     "type": "PATH",
  8.                     "value": "/usr/local/Ascend/latest"
  9.                 },
  10.                 ……
  11.         }
  12.     ]
  13. }
复制代码
在算子工程AddCustom目录下执行如下命令，举行算子工程编译。
  1. ./build.sh
复制代码
编译乐成后，会在当前目录下创建build_out目录，并在build_out目录下天生自界说算子安装包custom_opp__.run，例如“custom_opp_ubuntu_x86_64.run”。
4.2 自界说算子安装包部署

在自界说算子包所在路径下，执行如下命令，安装自界说算子包。
  1. ./custom_opp_<target os>_<target architecture>.run
复制代码
命令执行乐成后，自界说算子包中的相关文件将部署至当前情况的OPP算子库的vendors/customize目录中.
假如用户部署多个自界说算子包，可通过如下命令指定路径安装：
  1. ./custom_opp_<target os>_<target architecture>.run
  2. --install-path=<path>
复制代码
说明：假如部署算子包时通过设置–install-path参数指定了算子包的安装目录，则在利用自界说算子前，需要执行source       /vendors/<vendor_name>/bin/set_env.bash   命令，set_env.bash脚本中将自界说算子包的安装路径追加到情况变量ASCEND_CUSTOM_OPP_PATH中，使自界说算子在当前情况中生效。
4.3 查看部署后的目录结构

  1. ├── opp    // 算子库目录
  2. │   ├── built-in     // 内置算子所在目录
  3. │   ├── vendors     // 自定义算子所在目录
  4. │       ├── config.ini
  5. │       └── vendor_name1   // 自定义算子所在目录，若不指定路径安装，默认为“customize”
  6. │           ├── framework     //自定义算子插件库
  7. │           ├── op_impl
  8. │           │   └── ai_core
  9. │           │       └── tbe
  10. │           │           ├── config
  11. │           │           │   └── ${soc_version}     //昇腾AI处理器类型
  12. │           │           │       └── aic-${soc_version}-ops-info.json     //自定义算子信息库文件
  13. │           │           ├── vendor_name1_impl    //自定义算子实现代码文件
  14. │           │           │   └── dynamic
  15. │           │           │       ├── xx.cpp
  16. │           │           │       └── xx.py
  17. │           │           ├── kernel     //自定义算子二进制文件
  18. │           │           │   └── ${soc_version}     //昇腾AI处理器类型
  19. │           │           │   └── config
  20. │           │           └── op_tiling
  21. │           │               ├── lib
  22. │           │               └── liboptiling.so
  23. │           └── op_proto     //自定义算子原型库所在目录
  24. │               ├── inc
  25. │               │   └── op_proto.h
  26. │               └── lib
  27. │       ├── vendor_name2   // 存储厂商vendor_name2部署的自定义算子
  28. vendor_name1   // 自定义算子所在目录，若不指定路径安装，默认为“customize”
  29. vendor_name2   // 存储厂商vendor_name2部署的自定义算子
复制代码
5 算子ST测试

CANN开发套件包中提供了ST测试工具“msOpST”，用于天生算子的ST测试用例并在硬件情况中执行。
本节仅以AddCustom算子为例，介绍ST测试工具的关键执行流程。
5.1 编写测试界说文件AddCustom_case.json

创建算子ST测试用例界说文件“AddCustom_case.json”，例如存储到跟算子工程目录“AddCustom”同级别的“AddCustom_st”路径下。
“AddCustom_case.json”文件的样例如下，开发者可基于此文件定制修改。
  1. [
  2.     {
  3.         "case_name": "Test_AddCustom_001",
  4.         "op": "AddCustom",
  5.         "input_desc": [
  6.             {
  7.                 "format": [
  8.                     "ND"
  9.                 ],
  10.                 "type": [
  11.                     "float16"
  12.                 ],
  13.                 "shape": [8,2048],
  14.                 "data_distribute": [
  15.                     "uniform"
  16.                 ],
  17.                 "value_range": [
  18.                     [
  19.                         0.1,
  20.                         1.0
  21.                     ]
  22.                 ],
  23.                 "name": "x"
  24.             },
  25.             {
  26.                 "format": [
  27.                     "ND"
  28.                 ],
  29.                 "type": [
  30.                     "float16"
  31.                 ],
  32.                 "shape": [8,2048],
  33.                 "data_distribute": [
  34.                     "uniform"
  35.                 ],
  36.                 "value_range": [
  37.                     [
  38.                         0.1,
  39.                         1.0
  40.                     ]
  41.                 ],
  42.                 "name": "y"
  43.             }
  44.         ],
  45.         "output_desc": [
  46.             {
  47.                 "format": [
  48.                     "ND"
  49.                 ],
  50.                 "type": [
  51.                     "float16"
  52.                 ],
  53.                 "shape": [8,2048],
  54.                 "name": "z"
  55.             }
  56.         ]
  57.     }
  58. ]
复制代码
5.2 设置ST测试用例执行时依靠的情况变量

${INSTALL_DIR}体现CANN软件安装目录，例如，/usr/local/Ascend/ascend-toolkit/latest。{arch-os}为运行情况的架构和操作体系，arch体现操作体系架构，os体现操作体系，例如x86_64-linux。
  1. export DDK_PATH=${INSTALL_DIR}
  2. export NPU_HOST_LIB=${INSTALL_DIR}/{arch-os}/devlib
复制代码
5.3 天生测试用例

进入msOpST工具所在目录，执行如下命令天生并执行测试用例。
step1：进入msOpST工具所在目录。
  1. cd $HOME/Ascend/ascend-toolkit/latest/python/site-packages/bin
复制代码
step2：天生测试用例文件并执行.
  1. ./msopst run -i $HOME/AddCustom_st/AddCustom_case.json -soc <soc_version> -out $HOME/AddCustom_st
复制代码


  • -i：算子测试用例界说文件（*.json）的路径，可设置为绝对路径或者相对路径。
  • -soc：昇腾AI处理器的型号，请根据现真相况举行替换。
  • -out：天生文件所在路径。
    此命令执行完成后，会输出类似如下打屏效果：
  1. ------------------------------------------------------------------------
  2. - test case count: 1
  3. - success count: 1
  4. - failed count: 0
  5. ------------------------------------------------------------------------
  6. 2023-08-28 20:20:40 (25058) - [INFO] Process finished!
  7. 2023-08-28 20:20:40 (25058) - [INFO] The st report saved in:  xxxx/AddCustom_st/20230828202015/st_report.json.
复制代码
也可以查看上述屏**显信息提示的“st_report.json”**文件，查看详细运行效果。
参考学习：
基于自界说算子工程的算子开发-快速入门-Ascend C算子开发-算子开发-CANN社区版8.0.RC3.alpha003开发文档-昇腾社区

免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。

本帖子中包含更多资源

您需要 登录 才可以下载或查看，没有账号？立即注册

x
回复

使用道具 举报

0 个回复

倒序浏览
返回列表

快速回复

高级模式
B Color Image Link Quote Code Smilies
您需要登录后才可以回帖 登录 or 立即注册

本版积分规则

星球的眼睛

金牌会员
这个人很懒什么都没写!

楼主热帖

标签云

挺好的 服务器
微信订阅号
微信服务号
微信客服
小程序
H5
关于我们 商务合作 网站地图
©Copyright 2022-2024 杭州祥升科技有限公司 版权所有 浙ICP备20004199号
快速回复 返回顶部 返回列表