【2024第一期CANN训练营】Ascend C算子开发进阶篇

最新推荐文章于 2024-10-10 08:33:15 发布

小超编程

最新推荐文章于 2024-10-10 08:33:15 发布

阅读量345

点赞数 5

文章标签： php 数据库开发语言

本文链接：https://blog.csdn.net/2302_79615274/article/details/139204715

版权

相比基础篇的算子开发，实际中的算子开发场景会更为灵活和复杂：算子的输入shape等信息不是固定不变的，开发者需要根据此信息来决定数据的并行切分策略，也就是需要写Tiling算法；算子开发完成后要完成单算子调用和网络中调用，不再局限于<<<>>>的基础调用。

本篇内容将会介绍标准的算子开发流程，完成一个实际场景下的算子开发，包括工程创建、算子核函数实现、Host侧算子实现、算子工程编译部署以及单算子调用。我们将以一个简单的Add算子为例，定义算子类型为AddCustom，以展示动态shape的算子开发过程。

1. 工程创建

首先，我们需要创建一个新的Ascend C算子开发工程。CANN软件包中提供了工程创建工具msopgen，可以基于算子原型定义文件生成工程。

编写AddCustom算子的原型定义JSON文件。例如，创建一个名为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"]           }       ]   } ]

使用msopgen工具生成AddCustom算子的开发工程：

${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

替换<soc_version>为你的昇腾AI处理器型号，例如Ascend910A。

依次对以下5个文件进行开发：add_custom.cpp ，add_custom_tiling.h ，add_custom.cpp，CMakePresets.json，gen_data.py

AddCustom ├── build.sh          // 编译入口脚本 ├── cmake │   ├── config.cmake │   ├── util          // 算子工程编译所需脚本及公共编译文件存放目录 ├── CMakeLists.txt    // 算子工程的CMakeLists.txt ├── CMakePresets.json // 4.编译配置项 ├── framework         // 算子插件实现文件目录，单算子模型文件的生成不依赖算子适配插件，无需关注 ├── op_host                    // host侧实现文件 │   ├── add_custom_tiling.h    // 2.算子tiling定义文件 │   ├── add_custom.cpp         // 3.算子原型注册、shape推导、信息库、tiling实现等内容文件 │   ├── CMakeLists.txt ├── op_kernel                  // kernel侧实现文件 │   ├── CMakeLists.txt   │   ├── add_custom.cpp         // 1.算子核函数实现文件 ├── scripts                    // 自定义算子工程打包相关脚本所在目录 │   ├── gen_data.py            // 5.输入数据和真值数据生成脚本文件 │   ├── verify_result.py       // 验证输出数据和真值数据是否一致的验证脚本

2. Kernel侧核函数实现

2.1 核函数定义（add_custom.cpp）

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(tilingData, tiling);    // 初始化算子类    KernelAdd op;    // 算子类的初始化函数    op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);    if (TILING_KEY_IS(1)) {        // 完成算子实现的核心逻辑        op.Process();   } }

2.2 KernelAdd类实现

定义KernelAdd类，包含初始化、核心处理和流水任务实现。

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计算得到singleCoreSize(每个核上总计算数据大小)、tileNum（每个核上分块个数）、singleTileLength（每个分块大小）等变量        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;                // 获取当前核的起始索引        xGm.SetGlobalBuffer((__gm__ DTYPE_X*)x + this->blockLength * GetBlockIdx(), this->blockLength);        yGm.SetGlobalBuffer((__gm__ DTYPE_Y*)y + this->blockLength * GetBlockIdx(), this->blockLength);        zGm.SetGlobalBuffer((__gm__ DTYPE_Z*)z + this->blockLength * 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        LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();        LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>();         // 将GlobalTensor数据拷贝到LocalTensor        DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);        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从队列中取出，用于后续计算        LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();        LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>();        // 从Queue中分配输出Tensor        LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>();        // 调用Add接口进行计算        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        LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();        // 将输出Tensor拷贝到GlobalTensor中        DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);        // 将不再使用的LocalTensor释放        outQueueZ.FreeTensor(zLocal);   }private:    //Pipe内存管理对象    TPipe pipe;    //输入数据Queue队列管理对象，QuePosition为VECIN    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;    //输出数据Queue队列管理对象，QuePosition为VECOUT    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;    //管理输入输出Global Memory内存地址的对象，其中xGm, yGm为输入，zGm为输出    GlobalTensor<DTYPE_X> xGm;    GlobalTensor<DTYPE_Y> yGm;    GlobalTensor<DTYPE_Z> zGm;    // 每个核上总计算数据大小    uint32_t blockLength;    // 每个核上总计算数据分块个数    uint32_t tileNum;    // 每个分块大小    uint32_t tileLength; };

3. Host侧算子实现（add_custom_tiling.h ，add_custom.cpp）

完成核函数开发后，需要在Host侧实现算子的Tiling和原型注册。

3.1 Tiling定义（add_custom_tiling.h ）

修改add_custom_tiling.h文件，定义Tiling参数。

#ifndef ADD_CUSTOM_TILING_H #define ADD_CUSTOM_TILING_H #include "register/tilingdata_base.h"namespace optiling { 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类到对应的AddCustom算子 REGISTER_TILING_DATA_CLASS(AddCustom, TilingData) } #endif

3.2 Tiling实现（add_custom.cpp）

大多数情况下，Local Memory的存储，无法完整的容纳算子的输入与输出，需要每次搬运一部分输入进行计算然后搬出，再搬运下一部分输入进行计算，直到得到完整的最终结果，这个数据切分、分块计算的过程称之为Tiling。根据算子的shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循环多少次）的计算程序，称之为Tiling实现。

TilingData、block_dim、TilingKey、workspace这些概念的具体解释如下：

TilingData：切分算法相关参数，比如每次搬运的块大小，以及总共循环多少次，通过结构体存储，由开发者自行设计。
block_dim：算子数据切分的份数。例如，需要计算8M的数据，每个核上计算1M的数据，block_dim设置为8，但是为了充分利用硬件资源，一般将block_dim设置为硬件平台的核数，根据核数进行数据切分。
TilingKey（可选）：不同的kernel实现分支可以通过TilingKey来标识，host侧设置TilingKey后，可以选择对应的分支。例如，一个算子在不同的shape下，有不同的算法逻辑，kernel侧可以通过TilingKey来选择不同的算法逻辑，在host侧Tiling算法也有差异，host/kernel侧通过相同的TilingKey进行关联。
workspace size（可选）：workspace是设备侧Global Memory上的一块内存。在Tiling函数中可以设置workspace的大小，框架侧会为其在申请对应大小的设备侧Global Memory，在对应的算子kernel侧实现时可以使用这块workspace内存。

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->GetInputTensor(0)->GetShapeSize();                // 设置每个块的维度，设置TilingData        context->SetBlockDim(BLOCK_DIM);        tiling.set_totalLength(totalLength); // 设置总计算数据量        tiling.set_tileNum(TILE_NUM); // 设置每个核上的tile数量                // 将TilingData实例序列化并保存到TilingContext中，以便后续在kernel侧使用。        tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());        context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());                // 设置TilingKey（可选），用于选择kernel实现分支        context->SetTilingKey(1);                // 设置workspace大小（可选）        // 如果需要在设备侧Global Memory上分配workspace内存，可以通过GetWorkspaceSizes获取大小指针并设置。        size_t *currentWorkspace = context->GetWorkspaceSizes(1);        currentWorkspace[0] = 0;                return ge::GRAPH_SUCCESS;   } }

3.3 Shape推导函数实现（add_custom.cpp）

网络模型中的Shape推导和dtype推导

在实际的网络模型生成过程中，除了算子的调用，还需要进行Tensor的shape和数据类型（dtype）的推导。
推导的目的是为了在图执行之前确定各个Tensor的数据类型和形状，这样可以提前校验其正确性，并为算子的输出张量描述提供信息，包括形状、数据类型和数据排布格式。
通过提前推理出算子的输出张量描述，可以在算子构图准备阶段为所有张量静态分配内存，从而避免动态内存分配带来的性能开销。

AddCustom算子的InferShape实现

输出Tensor z的描述信息与输入Tensor x,y的描述信息相同，因此可以直接将任意一个输入Tensor的描述赋给输出Tensor。
InferShape函数的实现代码如下：

namespace ge {    static graphStatus InferShape(gert::InferShapeContext* context) {        const auto inputShape = context->GetInputShape(0); // 获取输入Tensor的形状        auto outputShape = context->GetOutputShape(0);     // 获取输出Tensor的形状        *outputShape = *inputShape;                        // 将输入形状赋给输出形状        return GRAPH_SUCCESS;                              // 返回成功状态   } }

3.4 算子原型注册（add_custom.cpp）

算子原型注册主要描述了算子的输入输出、属性等信息，以及算子在AI处理器上的相关实现信息。它还关联了Tiling实现、Shape推导等函数。这个过程对于确保算子能够正确地处理数据并在AI处理器上高效执行至关重要。

算子原型定义包括了算子的输入输出描述、数据类型、格式和属性等信息。例如，Add算子的输入x的描述信息包括它是必选的（ParamType REQUIRED），支持的数据类型（DataType），以及支持的格式（Format）。

在add_custom.cpp文件中注册算子原型。

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("ascend910");
        }
    };
    OP_ADD(AddCustom);
}

4. 算子编译部署（CMakePresets.json）

编译AddCustom工程，生成自定义算子安装包，并将其安装到算子库中。

修改CMakePresets.json中的ASCEND_CANN_PACKAGE_PATH为CANN软件包安装路径。

"ASCEND_CANN_PACKAGE_PATH": {
    "type": "PATH",
    "value": "/usr/local/Ascend/latest"  // 替换为CANN软件包安装后的实际路径
},
"CMAKE_CROSS_PLATFORM_COMPILER": {       // 替换为交叉编译工具安装后的实际路径
    "type": "PATH",
    "value": "/usr/bin/aarch64-linux-gnu-g++"
}

在算子工程目录下执行./build.sh命令进行编译。成功后，将在build_out目录下生成自定义算子安装包。

./build.sh

安装自定义算子包到算子

cd build_out
./custom_opp_euleros_aarch64.run

5. 单算子调用（AclNNInvocation）

5.1 准备验证代码工程

代码工程目录结构如下，您可以单击LINK，获取样例工程的完整样例：

依次修改op_runner.cpp，CMakeLists.txt，gen_data.py即可

├──input                         // 存放脚本生成的输入数据目录
├──output                        // 存放算子运行输出数据和真值数据的目录
├── inc                          // 头文件目录 
│   ├── common.h                 // 声明公共方法类，用于读取二进制文件 
│   ├── operator_desc.h          // 算子描述声明文件，包含算子输入/输出，算子类型以及输入描述与输出描述 
│   ├── op_runner.h              // 算子运行相关信息声明文件，包含算子输入/输出个数，输入/输出大小等 
├── src 
│   ├── CMakeLists.txt     // 2.编译规则文件
│   ├── common.cpp         // 公共函数，读取二进制文件函数的实现文件
│   ├── main.cpp    	   // 将单算子编译为om文件并加载om文件执行
│   ├── operator_desc.cpp  // 构造算子的输入与输出描述 
│   ├── op_runner.cpp      // 1.单算子编译与运行函数实现文件
├── scripts
│   ├── verify_result.py   // 真值对比文件
│   ├── gen_data.py        // 3.输入数据和真值数据生成脚本文件
│   ├── acl.json           // acl配置文件

5.2 单算子调用流程（op_runner.cpp）

头文件导入：编写单算子的调用代码时，要包含自动生成的单算子API执行接口头文件。

#include "aclnn_add_custom.h"

AscendCL初始化：执行此步骤以初始化AscendCL环境，为后续的算子调用和资源管理做准备。
运行管理资源申请：在此阶段，您需要设置设备ID并获取当前运行模式，这将影响后续接口调用流程（如数据传输等）。
申请内存存放算子的输入输出：为算子的输入和输出数据申请内存空间，确保算子可以正常读写数据。
传输数据：将主机内存中的数据传输到设备内存，以便算子可以在设备上进行计算。
计算workspace大小并申请内存：通过调用算子API的GetWorkspaceSize接口来确定执行算子所需的workspace大小，并据此申请相应的内存。

size_t workspaceSize = 0;
aclOpExecutor *handle = nullptr;
auto ret = aclnnAddCustomGetWorkspaceSize(inputTensor_[0], inputTensor_[1], outputTensor_[0],
                                          &workspaceSize, &handle);

执行算子：调用算子API的执行接口，开始算子的实际计算过程。

if (aclnnAddCustom(workspace, workspaceSize, handle, stream) != ACL_SUCCESS) {
    (void)aclrtDestroyStream(stream);
    ERROR_LOG("Execute Operator failed. error code is %d", static_cast<int32_t>(ret));
    return false;
}

同步等待：在计算完成后，使用同步操作等待所有计算任务完成，确保数据的完整性和准确性。
处理执行算子后的输出数据：根据应用需求，对算子的输出数据进行后续处理，如显示、存储或其他分析。
释放运行管理资源：在算子执行完毕后，释放之前申请的设备资源，维护系统的稳定性和性能。
AscendCL去初始化：最后，结束AscendCL环境，确保资源得到正确释放，避免潜在的资源泄露问题。

5.4 CMakeLists文件配置（CMakeLists.txt）

设置AUTO_GEN_PATH变量：确保该变量正确指向算子工程的build_out/autogen目录，这是为了CMake能够找到自动生成的单算子API执行接口头文件和源文件。

set(AUTO_GEN_PATH "../../AddCustom/build_out/autogen")

增加头文件搜索路径：在CMakeLists.txt中，添加include_directories指令，将AUTO_GEN_PATH变量所指向的目录加入到头文件的搜索路径中，以便编译器能够找到并包含这些自动生成的头文件。

include_directories(
    ${AUTO_GEN_PATH}
)

生成可执行文件规则：在CMakeLists.txt中，使用add_executable指令来定义可执行文件的生成规则，包括目标名称和源文件列表。确保将自动生成的单算子API调用实现文件（如aclnn_add_custom.cpp）包含在内，以便链接成最终的可执行文件。

add_executable(execute_add_op
    ${AUTO_GEN_PATH}/aclnn_add_custom.cpp
)

链接nnopbase库：在CMakeLists.txt中，通过target_link_libraries指令将nnopbase库链接到您的可执行文件目标上。这是为了确保在运行时，程序能够正确链接到所需的库文件，从而能够调用单算子API。

target_link_libraries(execute_add_op
    nnopbase
)

5.5 数据生成（gen_data.py ）

以固定shape的add_custom算子为例，输入数据和真值数据生成的脚本样例如下：根据算子的输入输出编写脚本，生成输入数据和真值数据。

#!/usr/bin/python3 
# -*- coding:utf-8__  
# 版权所有 (c) 华为技术有限公司 2022-2023。

import numpy as np  # 导入numpy库，用于科学计算

# 定义生成基准数据的函数
def gen_golden_data_simple():
    # 生成两个大小为[8, 2048]的随机数矩阵，范围在[-100, 100]之间，数据类型为float16
    input_x = np.random.uniform(-100, 100, [8, 2048]).astype(np.float16)
    input_y = np.random.uniform(-100, 100, [8, 2048]).astype(np.float16)
    
    # 计算两个矩阵的和，得到基准数据（golden data），数据类型为float16
    golden = (input_x + input_y).astype(np.float16)
    
    # 将生成的输入矩阵input_x和input_y分别保存到二进制文件"./input/input_x.bin"和"./input/input_y.bin"
    # 将基准数据golden保存到二进制文件"./output/golden.bin"，用于后续结果验证
    input_x.tofile("./input/input_x.bin")
    input_y.tofile("./input/input_y.bin")
    golden.tofile("./output/golden.bin")

if __name__ == "__main__":
    # 调用函数生成基准数据
    gen_golden_data_simple()

5.6 编译与运行（run.sh）

cd AclNNInvocation
    
// 看到test pass则精度验证通过
bash run.sh

小超编程

关注

5
点赞
踩
6

收藏

觉得还不错? 一键收藏
打赏
0
评论
复制链接

分享到 QQ

分享到新浪微博

扫一扫