CANN学习资源开源仓的算子贡献开发一

第04部分是cube算子，第05部分是融合算子，暂时略过，下面是06开源仓贡献部分。下载源码编译安装gawk5.3.0因为镜像无此命令。开源算子仓提供四类仓库，是内置算子的集合，支持多个算子同时编译出包。

以ops-math为例，仓库出包有三种方式：

• 自定义算子包：选择部分算子编译生成的包，以挂载形式作用于CANN包，不改变原始包内容。优先级高于原始CANN包。该包支持aclnn和图模式调用AI Core、AI CPU算子。
• ops-math包：选择整个项目编译生成的包称为ops-math包，可完整替换CANN包对应部分。该包支持aclnn和图模式调用AI Core算子。
• ops-math静态库：整个项目编译为一个静态库文件，包含libcann_math_static.a和aclnn接口头文件。仅支持aclnn调用AI Core算子。

然后以自定义算子包方式进行experimental/math/arange算子的完整验证，其中编译参数–soc：${soc_version}表示NPU型号。Atlas A2 训练/推理系列产品使用"ascend910b"（默认），Atlas A3 训练/推理系列产品使用"ascend910_93"，Ascend 950PR/Ascend 950DT产品使用"ascend950"。其中执行参数${mode}：表示执行模式，目前支持eager（aclnn调用）、graph（图模式调用）。${pkg_mode}：表示包模式，目前仅支持cust自定义算子包。

# 编译、部署、执行
bash build.sh --pkg --experimental --vendor_name=custom --soc=ascend910b --ops=arange -j16

./build_out/cann-ops-math-custom_linux-aarch64.run --install-path=${HOME}/

source ${HOME}/vendors/custom_math/bin/set_env.bash;bash build.sh --experimental --run_example arange eager cust --vendor_name=custom

验证过后，再说开发。使用build.sh在指定目录下生成空的算子工程模板。同样想说的是，工程模版只有框架、样子齐全，但内容不足，如tiling函数文件太空了，不如拷贝别人实现好的工程改吧改吧更省事。

# bash build.sh --genop=${op_class}/${op_name}
bash build.sh --genop=experimental/math/add_custom

Tiling 实现需交付三个：${op_name}_tiling.cpp、${op_name}_tiling_key.h、${op_name}_tiling_data.h（Tiling 结构体定义文件）。由于 TilingKey 通常表现为无明确语义的长数字串，易记易理解性较差，因此开源仓设计了 Tiling 模板编程方法来简化使用，并要求使用。其中，${op_name}_tiling_key.h 用于定义 TilingKey 模板参数，如下：

/*!
 * \file add_custom_tiling_key.h
 * \brief add_custom tiling key declare
 */

#ifndef __ADD_CUSTOM_TILING_KEY_H__
#define __ADD_CUSTOM_TILING_KEY_H__

#include "ascendc/host_api/tiling/template_argument.h"

/* Mode场景定义 */
#define ELEMENTWISE_TPL_SCH_MODE_0 0
#define ELEMENTWISE_TPL_SCH_MODE_1 1
/* 继续定义其他Mode场景... */

/* 定义了模板参数 schMode，其取值范围限定为 0 或 1*/
ASCENDC_TPL_ARGS_DECL(
    AddCustom,
    ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));

/* 模板参数组合 当前仅配置了一组模板参数（即 schMode 可在 0、1 中任选其一）*/
ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
    ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));

#endif

//因为该算子无需配置多组 TilingKey，简化修改如下：仅保留 schMode 取值为 0 的配置
#define ELEMENTWISE_TPL_SCH_MODE_0 0

ASCENDC_TPL_ARGS_DECL(
    AddCustom,
    ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0));

ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
    ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0)));

TilingKey 的生成逻辑则在 ${op_name}_tiling.cpp 中实现：

#include "log/log.h"
#include "util/math_util.h"
#include "op_host/tiling_util.h"
#include "op_host/tiling_templates_registry.h"
#include "../op_kernel/add_custom_tiling_data.h"
#include "../op_kernel/add_custom_tiling_key.h"

namespace optiling {

static ge::graphStatus TilingFunc(gert::TilingContext *context)
{
    uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();
    context->SetBlockDim(8);
    AddCustomTilingData *tiling = context->GetTilingData<AddCustomTilingData>();
    tiling->totalLength = totalLength;
    tiling->tileNum = 1;
    //tiling模板化编程
    uint64_t tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0);
    context->SetTilingKey(tilingKey);

    return ge::GRAPH_SUCCESS;
}
//本文件仅包含 TilingFunc 函数，新增 IMPL_OP_OPTILING 宏定义，这是开源仓算子与自定义算子工程最核心的区别
IMPL_OP_OPTILING(AddCustom).Tiling(TilingFunc);
}  // namespace optiling

Kernel 实现需交付两个：op_name.cpp（入口文件）和op_name.h（实现文件）。也可和自定义算子一样，入口和实现均写在${op_name}.cpp文件中。入口文件如下：

#include "add_custom.h"

enum class AddTilingKey : uint32_t
{
    TILING_KEY_EXAMPLE = 0,
};

template <uint32_t schMode>
 __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{
    REGISTER_TILING_DEFAULT(AddCustomTilingData);
    GET_TILING_DATA_WITH_STRUCT(AddCustomTilingData, tiling_data, tiling);
    //核函数的标准编写范式：尽管 schMode 仅支持取值 0，此处的条件判断看似无实际意义，但该写法为算子预留了可扩展空间
    if constexpr (schMode == static_cast<uint32_t>(AddTilingKey::TILING_KEY_EXAMPLE)) {
        KernelAdd<DTYPE_X1, DTYPE_X2, DTYPE_Y> op;
        op.Init(x, y, z, tiling_data.totalLength, tiling_data.tileNum);
        op.Process();
    }
}

这里有一点觉得累赘的就是要定义enum class AddTilingKey，为什么不能用${op_name}_tiling_key.h里面已定义的值，而要重复定义本是同一个东西的另一个东西呢？

实现文件如下：

#include "kernel_operator.h"
#include "kernel_tiling/kernel_tiling.h"
#include "add_custom_tiling_data.h"
#include "add_custom_tiling_key.h"

constexpr int32_t BUFFER_NUM = 1;  // tensor num for each queue

template <class dtypeX, class dtypeY, class dtypeZ>
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)
    {}
//略    
}

如果仅用 aclnn 调用单算子（不入图），就不用写 Shape/Dtype 推导函数。否则，当一个算子被加入计算图时，为确保图的正确性和编译、优化、执行流程顺利进行，通常需要为该算子实现推导函数：

• InferShape：用于推导输出张量的形状。
• InferDataType：用于推导输出张量的数据类型。

这样可以在图执行之前，知道各Tensor的数据类型和形状，提前校验其正确性；同时提前推理出算子的输出张量描述，包括张量的形状、数据类型及数据排布格式等信息，算子构图准备阶段就可以为所有的张量静态分配内存，避免动态内存分配带来的开销。

形状推导在 op_host/add_custom_infershape.cpp

数据类型推导在 op_graph/add_custom_graph_infer.cpp 我也是不明白，为什么要分开2个不同的目录

图模式调用需要将算子原型注册到Graph Engine（简称GE）中，以便GE能够识别该类算子的输入、输出及属性信息。注册通过REG_OP接口完成。文件是 op_graph/add_custom_proto.h。

再后面的编译部署调用，和前面验证的类似，就不写了。最后，基于 Ascend C 开展算子工程化开发时，开发者可选择自定义算子工程或开源算子仓工程两种形式：