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

最新推荐文章于 2024-12-03 23:14:49 发布

绿洲213

最新推荐文章于 2024-12-03 23:14:49 发布

阅读量2.1k

点赞数 29

分类专栏： 2024·CANN训练营第一季文章标签： c语言 php 开发语言

本文链接：https://blog.youkuaiyun.com/weixin_46227276/article/details/136806866

版权

2024·CANN训练营第一季专栏收录该内容

6 篇文章

订阅专栏

文章目录

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

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

相比基础篇的算子开发，实际中的算子开发场景会更为灵活和复杂：算子的输入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