1 融合算子概述

融合算子是将多个独立的基础算子整合为单一整体算子的高性能计算单元，其计算功能与多算子串行执行的计算逻辑完全等价。通过算子融合技术，可消除多算子迭代执行过程中冗余的数据搬移与缓存读写开销，充分挖掘昇腾AI芯片的硬件并行能力，显著提升深度学习模型的推理与训练效率。

依据硬件计算单元的协同工作类型，昇腾架构下的融合算子可分为两类：一是VV融合算子，即多个Vector单元算子的融合；二是CV融合算子（Mix融合算子），即Cube矩阵计算单元与Vector通用计算单元的跨单元融合，也是高性能矩阵运算优化的核心方案。

2 融合算子性能优化原理

2.1 VV融合性能优化机制

传统多Vector算子独立执行时，每一个算子完成计算后均需将片上缓存数据搬移至全局内存，下一算子执行前再重新读取数据，频繁的内存交互会产生大量冗余耗时。VV融合通过合并连续Vector算子的计算流程，仅保留一次输入数据搬移与一次输出数据落盘操作，省略中间数据读写环节，大幅降低数据传输开销，提升计算流水线利用率。

为直观验证VV融合的优化效果，通过模拟时序实验对比融合前后的执行耗时，实验仿真代码如下：

import time
def mock_data_move(data, direction):
    time.sleep(0.2)
    return data

def mock_op1_compute(data):
    time.sleep(0.05)
    return [x + 1 for x in data]

def mock_op2_compute(data):
    time.sleep(0.05)
    return [x * 2 for x in data]

print("===== 非融合算子执行流程 =====")
start = time.time()
data_in1 = mock_data_move([1,2,3], "数据搬入设备")
res1 = mock_op1_compute(data_in1)
data_out1 = mock_data_move(res1, "数据搬出设备")
data_in2 = mock_data_move(data_out1, "数据搬入设备")
res2 = mock_op2_compute(data_in2)
final_non_fused = mock_data_move(res2, "数据搬出设备")
non_fused_time = time.time() - start
print(f"非融合总耗时：{non_fused_time:.2f}秒 | 结果：{final_non_fused}\n")

print("===== 融合算子执行流程 =====")
start = time.time()
data_in = mock_data_move([1,2,3], "数据搬入设备")
res1 = mock_op1_compute(data_in)
res2 = mock_op2_compute(res1)
final_fused = mock_data_move(res2, "数据搬出设备")
fused_time = time.time() - start
print(f"融合总耗时：{fused_time:.2f}秒 | 结果：{final_fused}\n")

speed_up = (non_fused_time - fused_time) / non_fused_time * 100
print(f"融合算子耗时减少{speed_up:.0f}%！")

仿真结果表明，VV融合通过精简数据交互流程，可有效降低算子迭代执行的总耗时，实现计算效率的显著提升。

2.2 CV融合性能优化机制

CV融合针对Cube矩阵计算与Vector通用计算的串行场景设计，传统执行模式下，Cube单元完成矩阵运算后需将结果写回全局内存，Vector单元读取数据后再执行激活、偏置叠加等后续运算，串行执行模式存在严重的流水线等待与数据搬移冗余。

CV融合依托昇腾A2分离架构，实现Cube与Vector单元的流水线并行计算，Cube输出结果可直接送入Vector输入缓存（L0C），无需落地统一缓存（UB）与全局内存，实现计算链路的无缝衔接。通过分块流水线并行策略，实现数据加载、矩阵计算、通用运算、数据落盘的并行重叠执行，最大化挖掘硬件算力。

为量化CV融合的优化收益，构建矩阵计算+通用运算的时序仿真模型，设定硬件各阶段固定耗时，对比融合前后的整体执行时延，仿真代码如下：

import numpy as np

# 硬件各阶段耗时定义（单位：秒）
CUBE_MOVE_IN = 0.1     # Cube数据搬入耗时
CUBE_COMPUTE = 0.2     # Cube矩阵运算耗时
CUBE_MOVE_OUT = 0.1    # Cube数据搬出耗时
VECTOR_MOVE_IN = 0.05  # Vector数据搬入耗时
VECTOR_COMPUTE = 0.1   # Vector通用运算耗时
VECTOR_MOVE_OUT = 0.05 # Vector数据搬出耗时

# 单模块完整耗时
SINGLE_CUBE = CUBE_MOVE_IN + CUBE_COMPUTE + CUBE_MOVE_OUT  # 0.4s
SINGLE_VECTOR = VECTOR_MOVE_IN + VECTOR_COMPUTE + VECTOR_MOVE_OUT  # 0.2s

# 融合前：Cube全量计算串行Vector全量计算
def matmul_add_original():
    block_count = 4
    total_cube = block_count * SINGLE_CUBE  # 4块Cube串行计算
    total_vector = block_count * SINGLE_VECTOR  # Cube完成后Vector串行计算
    return total_cube + total_vector

# 融合后：多块数据流水线并行执行
def matmul_add_fused():
    # 首块数据Cube独立计算
    step1 = SINGLE_CUBE
    # 后续多组Cube与Vector流水线并行
    step2 = max(VECTOR_MOVE_IN, CUBE_MOVE_IN)
    step3 = max(VECTOR_COMPUTE, CUBE_COMPUTE)
    step4 = max(VECTOR_MOVE_OUT, CUBE_MOVE_OUT)
    step5 = max(VECTOR_MOVE_IN, CUBE_MOVE_IN)
    step6 = max(VECTOR_COMPUTE, CUBE_COMPUTE)
    step7 = max(VECTOR_MOVE_OUT, CUBE_MOVE_OUT)
    step8 = max(VECTOR_MOVE_IN, CUBE_MOVE_IN)
    step9 = max(VECTOR_COMPUTE, CUBE_COMPUTE)
    step10 = max(VECTOR_MOVE_OUT, CUBE_MOVE_OUT)
    # 末块数据Vector独立计算
    step11 = VECTOR_MOVE_IN
    step12 = VECTOR_COMPUTE
    step13 = VECTOR_MOVE_OUT
    total_time = step1 + step2 + step3 + step4 + step5 + step6 + step7 + step8 + step9 + step10 + step11 + step12 + step13
    return total_time

# 性能对比测试
if __name__ == "__main__":
    time_original = matmul_add_original()
    time_fused = matmul_add_fused()
    speedup = (time_original - time_fused) / time_original * 100
    print("=== 昇腾Matmul+Add融合算子性能测试 ===")
    print(f"融合前逻辑：全量Cube矩阵计算→全量Vector通用计算（纯串行）")
    print(f"融合前总耗时：{time_original:.2f} 秒")
    print(f"\n融合后逻辑：分块流水线并行计算（Cube/Vector异步协同）")
    print(f"融合后总耗时：{time_fused:.2f} 秒")
    print(f"\n性能提升比例：{speedup:.1f}%")

仿真结果验证，CV融合通过消除串行等待、重叠硬件操作，可大幅降低整体计算时延，实现深度学习矩阵运算场景的算力提速。

3 VV融合算子完整设计与实现

VV融合算子开发分为Host侧与Kernel侧两大模块：Host侧运行于CPU，负责算子编译配置、形状推导、数据类型推导与分片参数下发；Kernel侧运行于AI Core，负责硬件层数据搬移与核心计算逻辑。本节以平方差融合算子为例，完成标准化开发实现。

3.1 Host侧开发实现

3.1.1 分片参数结构体定义

自定义分片结构体，用于存储算子运行所需的全局数据长度、分块数量等自定义参数，实现Host与Kernel侧的参数交互。

//square_diff_tiling.h
#ifndef SQUARE_DIFF_TILING_H
#define SQUARE_DIFF_TILING_H
#include<cstdint>

// VV融合算子自定义分片参数结构体
struct SquareDiffTilingData{
    uint32_t totalLength;  // 全局数据总长度
    uint32_t tilieNum;     // 单核分块迭代次数
};

#endif

3.1.2 算子注册与编译逻辑实现

Host侧核心逻辑包含算子参数注册、形状推导、数据类型推导与分片配置四部分，为GE图引擎提供算子编译、调度与内存分配依据。

//op_host/square_diff.cpp
#include "../op_kernel/square_diff_tiling.h"
#include "register/op_def_registry.h"

namespace optiling{
// 分片配置函数：完成分块参数计算、并行核数配置、工作空间申请
    static ge::graphStatus TilingFunc(gert::TilingContext *context){
        SquareDiffTilingData *tiling =context->GetTilingData<SquareDiffTilingData>();
        // 获取输入张量全局元素总数
        uint32_t totalLength=context->GetInputShape(0)->GetOriginShape().GetShapeSize();
        tiling->totalLength=totalLength;
        tiling->tileNum=1;
        // 配置8核并行调度
        context->SetBlockDim(8);
        // 无需用户自定义临时显存，工作空间大小置0
        size_t *currentWorkspace=context->GetWorkspaceSizes(1);
        currentWorkspace[0]=0;
        return ge::GRAPH_SUCCESS;
    }
}

namespace ge{
// 输出形状推导：VV逐元素运算，输出与输入维度一致
    static ge::graphStatus InferShape(gert::InferShapeContext* context){
        const gert::Shape* x1_shape=context->GetInputShape(0);
        gert::Shape* y_shape =context->GetOutputShape(0);
        *y_shape=*x1_shape;
        return GRAPH_SUCCESS;
    }

// 输出数据类型推导：继承输入数据类型
    static ge::graphStatus InferDataType(gert::InferDataTypeContext *context){
        const auto inputDataType =context->GetInputDataType(0);
        context->SetOutputDataType(0,inputDataType);
        return ge::GRAPH_SUCCESS;
    }
}

namespace ops{
// 算子注册类：定义算子输入输出属性、绑定编译与分片逻辑
    class SquareDiff:public OpDef{
    public:
        explicit SquareDiff(const char*name):OpDef(name){
            // 输入参数x配置
            this->Input("x")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16,ge::FORMAT_ND})
                .Format({ge::FORMAT_ND,ge::FORMAT_ND})
                .UnknownShapeFormat({ge::FORMAT_ND,ge::FORMAT_ND});
            // 输入参数y配置
            this->Input("y")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
                .Format({ge::FORMAT_ND, ge::FORMAT_ND})
                .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
            // 输出参数z配置
            this->Output("z")
                .ParamType(REQUIRED)
                .DataType({ge::DT_FLOAT16, ge::DT_FLOAT})
                .Format({ge::FORMAT_ND, ge::FORMAT_ND})
                .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
            // 绑定形状与类型推导函数
            this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);
            // 绑定分片函数与适配芯片型号
            this->AICore()
                .SetTiling(optiling::TilingFunc);
            this->AICore().AddConfig("ascend910b");
        }
    };
    // 注册自定义算子
    OP_ADD(SquareDiff);
}

3.2 Kernel侧开发实现

Kernel侧运行于AI Core，采用标准化的Init-Process-Copy架构，完成数据初始化、流水线循环、数据搬移与融合计算逻辑，核心计算模块实现加法、减法、乘法的VV融合运算。

//op_kernel/square_diff.cpp
#include "kernel_operator.h"
#include "square_diff_tiling.h"

constexpr int32_t BUFFER_NUM = 1;    // 队列缓冲区数量
constexpr int32_t QUEUE_DEPTH = 1;  // 队列深度

// 平方差融合算子内核类
class KernelSquareDiff {
public:
    __aicore__ inline KernelSquareDiff(){}
    // 初始化：配置全局内存、片上队列与分块参数
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
    {
        ascendc_assert(tileNum != 0, "tileNum can not be zero.\n");
        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.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));
    }

    // 主流程：循环执行数据搬入-计算-搬出
    __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:
    // 数据搬入：全局内存→片上缓存
    __aicore__ inline void CopyIn(int32_t progress)
    {
        AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.AllocTensor<DTYPE_Y>();
        AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
        AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }

    // 融合计算：(x+y)*(x-y) 平方差逻辑
    __aicore__ inline void Compute(int32_t progress)
    {
        AscendC::LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        AscendC::LocalTensor<DTYPE_Y> yLocal = inQueueY.DeQue<DTYPE_Y>();
        AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.AllocTensor<DTYPE_Z>();
        // VV融合：连续执行加法、减法、乘法运算
        AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);
        AscendC::Sub(xLocal, xLocal, yLocal, this->tileLength);
        AscendC::Mul(zLocal, zLocal, xLocal, this->tileLength);
        outQueueZ.EnQue<DTYPE_Z>(zLocal);
        // 释放片上缓存
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }

    // 数据搬出：片上缓存→全局内存
    __aicore__ inline void CopyOut(int32_t progress)
    {
        AscendC::LocalTensor<DTYPE_Z> zLocal = outQueueZ.DeQue<DTYPE_Z>();
        AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
        outQueueZ.FreeTensor(zLocal);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::TPosition::VECIN, QUEUE_DEPTH> inQueueX, inQueueY;
    AscendC::TQue<AscendC::TPosition::VECOUT, QUEUE_DEPTH> outQueueZ;
    AscendC::GlobalTensor<DTYPE_X> xGm;
    AscendC::GlobalTensor<DTYPE_Y> yGm;
    AscendC::GlobalTensor<DTYPE_Z> zGm;
    uint32_t blockLength;
    uint32_t tileNum;
    uint32_t tileLength;
};

// 算子内核入口函数
extern "C" __global__ __aicore__ void square_diff(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) {
    REGISTER_TILING_DEFAULT(SquareDiffTilingData);
    GET_TILING_DATA(tilingData, tiling);
    KernelSquareDiff op;
    op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);
    op.Process();
}

4 CV融合算子完整设计与实现

本文实现Matmul+Bias+LeakyReLU CV融合算子，依托昇腾A2分离模式，实现Cube矩阵乘、偏置叠加、Vector激活的跨单元流水线融合，规避中间数据落盘开销，适配大规模矩阵运算场景。

4.1 核心基础概念说明

4.1.1 TCubeTiling分片结构体

TCubeTiling是CANN框架预定义的矩阵乘分片结构体，分为Host侧与Device侧两个同名结构体，二者内存布局完全一致：Host侧optiling::TCubeTiling用于CPU端分片计算，Device侧AscendC::tiling::TCubeTiling用于AI Core端读取分片参数，封装了矩阵维度、分块尺寸、调度通路、迭代规则等全部矩阵乘所需分片信息。

4.1.2 硬件缓存约束

CV融合算子需严格适配硬件缓存规格：Cube输出缓存（L0C）容量为64KB，统一缓存（UB）为数百KB级别。矩阵分块尺寸baseM、baseN需满足：baseM * baseN * sizeof(float) ≤ 64KB，确保片上缓存存储合法，避免内存越界。

4.2 自定义分片结构体定义

拓展官方矩阵乘分片结构体，新增LeakyReLU斜率参数，实现融合算子超参与分片参数的统一下发。

//op_kernel/matmul_leakyrelu_custom_tiling.h
#ifndef MATMUL_LEAKYRELU_CUSTOM_TILING_H
#define MATMUL_LEAKYRELU_CUSTOM_TILING_H
#include <cstdint>
#include "kernel_tiling/kernel_tiling.h"

// CV融合算子自定义分片参数结构体
struct MatmulLeakyreluCustomTilingData {
    float alpha;                                  // LeakyReLU负区间斜率超参
	AscendC::tiling::TCubeTiling cubeTilingData;  // 官方矩阵乘分片参数
};

// 标准化宏定义写法（等价替代）
/*
BEGIN_TILING_DATA_DEF(MatmulLeakyreluCustomTilingData)
TILING_DATA_FIELD_DEF_STRUCT(optiling::TCubeTiling, cubeTilingData);
TILING_DATA_FIELD_DEF(float, alpha);
END_TILING_DATA_DEF;
*/
#endif

4.3 Host侧开发实现

Host侧完成矩阵维度配置、硬件分片策略生成、并行核数配置、工作空间申请、算子属性注册与张量形态推导。

//op_host/matmul_leakyrelu_custom.cpp
#include "../op_kernel/matmul_leakyrelu_custom_tiling.h"
#include "register/op_def_registry.h"
#include "tiling/platform/platform_ascendc.h"
#include "tiling/tiling_api.h"

using namespace matmul_tiling;

namespace optiling{
// 分片配置核心函数
static ge::graphStatus TilingFunc(gert::TilingContext* context){
	MatmulLeakyreluCustomTilingData *tiling =context ->GetTilingData<MatmulLeakyreluCustomTilingData>();
    // 固定矩阵运算维度
    int32_t M=1024;
    int32_t N=640;
    int32_t K=256;
    // 单步Cube计算分块尺寸
    int32_t baseM = 128;
    int32_t baseN = 128;

    // 初始化芯片平台与矩阵分片计算器
    auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
    MultiCoreMatmulTiling cubeTiling(ascendcPlatform);
    cubeTiling.SetDim(2);  // A2分离模式，固定2路AIV调度通路

    // 配置各张量存储位置、格式与数据类型
    cubeTiling.SetAType(TPosition::GM, CubeFormat::ND, DataType::DT_FLOAT16);
    cubeTiling.SetBType(TPosition::GM, CubeFormat::ND, DataType::DT_FLOAT16);
    cubeTiling.SetCType(TPosition::VECIN, CubeFormat::ND, DataType::DT_FLOAT);
    cubeTiling.SetBiasType(TPosition::GM, CubeFormat::ND, DataType::DT_FLOAT);

    // 配置矩阵形状与分块规则
    cubeTiling.SetShape(M, N, K);
    cubeTiling.SetOrgShape(M, N, K);
    cubeTiling.SetFixSplit(baseM, baseN, -1);
    cubeTiling.SetBias(true);
    cubeTiling.SetBufferSpace(-1, -1, -1);

    // 生成分片参数，校验合法性
    if (cubeTiling.GetTiling(tiling->cubeTilingData) == -1) { 
        return ge::GRAPH_FAILED;
    }

    // 配置LeakyReLU超参
    tiling->alpha = 0.001; 
    // 单核并行调度
    context->SetBlockDim(1);

    // 工作空间计算：系统库空间+用户自定义空间
    size_t userWorkspaceSize = 0;
    size_t systemWorkspaceSize = static_cast<size_t>(ascendcPlatform.GetLibApiWorkSpaceSize());
    size_t *currentWorkspace = context->GetWorkspaceSizes(1);
    currentWorkspace[0] = userWorkspaceSize + systemWorkspaceSize;

    return ge::GRAPH_SUCCESS;
}
}

namespace ge {
// 输出形状推导：遵循矩阵乘维度规则 [M,K]×[K,N]→[M,N]
static ge::graphStatus InferShape(gert::InferShapeContext* context)
{
    const gert::Shape* a_shape = context->GetInputShape(0);
    const gert::Shape* b_shape = context->GetInputShape(1);
    gert::Shape* out_shape = context->GetOutputShape(0);
    int64_t M = a_shape->GetDim(0);
    int64_t N = b_shape->GetDim(1);
    out_shape->SetDim(0, M);
    out_shape->SetDim(1, N);
    return ge::GRAPH_SUCCESS;
}

// 输出数据类型推导
static ge::graphStatus InferDataType(gert::InferDataTypeContext *context)
{
    const auto inputDataType = context->GetInputDataType(0);
    context->SetOutputDataType(0, inputDataType);
    return ge::GRAPH_SUCCESS;
}
}

namespace ops {
// CV融合算子注册类
class MatmulLeakyreluCustom : public OpDef {
public:
    explicit MatmulLeakyreluCustom(const char* name) : OpDef(name)
    {
        // 输入输出参数属性配置
        this->Input("a")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});
        this->Input("b")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});
        this->Input("bias")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});
        this->Output("c")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});

        // 绑定编译逻辑
        this->SetInferShape(ge::InferShape).SetInferDataType(ge::InferDataType);
        this->AICore()
            .SetTiling(optiling::TilingFunc);
        this->AICore().AddConfig("ascend910b");
    }
};
OP_ADD(MatmulLeakyreluCustom);
}

4.4 Kernel侧开发实现

Kernel侧实现矩阵乘、偏置叠加、LeakyReLU激活的流水线融合逻辑，包含多核偏移计算、片上队列管理、跨单元数据流转与异步数据搬移，完整复用VECIN直通优化机制，规避中间数据落地开销。

//op_kernel/matmul_leakyrelu_custom.cpp
#include "kernel_operator.h"
#include "matmul_leakyrelu_custom_tiling.h"
#include "lib/matmul_intf.h"
using namespace matmul;

// 向上取整工具函数
__aicore__ inline uint32_t Ceiling(uint32_t a, uint32_t b)
{
    return (a + b - 1) / b;
}

// CV融合算子内核模板类
template <typename aType, typename bType, typename cType, typename biasType> 
class MatmulLeakyKernel {
public:
    __aicore__ inline MatmulLeakyKernel(){};
    __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace,const TCubeTiling &tiling, float alpha, AscendC::TPipe *pipe);
    __aicore__ inline void Process(AscendC::TPipe *pipe);
    __aicore__ inline void MatmulCompute();
    __aicore__ inline void LeakyReluCompute();
    __aicore__ inline void CopyOut(uint32_t count);
    __aicore__ inline void CalcOffset(int32_t blockIdx, const TCubeTiling &tiling, int32_t &offsetA, int32_t &offsetB,int32_t &offsetC, int32_t &offsetBias);

    // 矩阵乘对象：配置GM输入、VECIN直通输出、ND数据格式
    Matmul<
    MatmulType<AscendC::TPosition::GM,CubeFormat::ND,aType>,
    MatmulType<AscendC::TPosition::GM, CubeFormat::ND,bType>,
    MatmulType<AscendC::TPosition::VECIN, CubeFormat::ND,cType>,
    MatmulType<AscendC::TPosition::GM, CubeFormat::ND,biasType>
	> matmulObj;

    // 全局内存张量绑定
    AscendC::GlobalTensor<aType> aGlobal;
    AscendC::GlobalTensor<bType> bGlobal;
    AscendC::GlobalTensor<cType> cGlobal;
    AscendC::GlobalTensor<biasType> biasGlobal;

    AscendC::LocalTensor<cType> reluOutLocal;
    float alpha;
    TCubeTiling tiling;
    AscendC::TQue<AscendC::QuePosition::VECOUT, 1>reluOutQueue_;
};

// 内核初始化：参数赋值、内存绑定、队列初始化、多核偏移适配
template <typename aType, typename bType, typename cType, typename biasType> 
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace,const TCubeTiling &tiling, float alpha, AscendC::TPipe *pipe)
{
    this->tiling = tiling;
    this->alpha = alpha;
    // 绑定全局内存张量维度
    aGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ aType *>(a), tiling.M * tiling.Ka);
    bGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ bType *>(b), tiling.Kb * tiling.N);
    cGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ cType *>(c), tiling.M * tiling.N);
    biasGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ biasType *>(bias), tiling.N);

    // 计算多核并行偏移量
    int offsetA = 0;
    int offsetB = 0;
    int offsetC = 0;
    int offsetBias = 0;
    CalcOffset(AscendC::GetBlockIdx(), tiling, offsetA, offsetB, offsetC, offsetBias); 

    // 绑定当前核负责的内存区间
    aGlobal = aGlobal[offsetA];
    bGlobal = bGlobal[offsetB];
    cGlobal = cGlobal[offsetC];
    biasGlobal = biasGlobal[offsetBias];

    // 初始化输出队列缓冲区
    pipe->InitBuffer(reluOutQueue_, 1, tiling.baseM * tiling.baseN * sizeof(cType)); 
}

// 主计算流程：矩阵迭代计算、融合运算、数据落盘
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::Process(AscendC::TPipe *pipe)
{
    uint32_t computeRound = 0;
    matmulObj.SetTensorA(aGlobal);
    matmulObj.SetTensorB(bGlobal);
    matmulObj.SetBias(biasGlobal);
    // 分块迭代完成全量矩阵计算
    while (matmulObj.template Iterate<true>()) {
        MatmulCompute();    // Cube矩阵乘计算
        LeakyReluCompute(); // Vector激活计算
        CopyOut(computeRound); // 数据异步落盘
        computeRound++;
    }
    matmulObj.End();
}

// 矩阵乘计算：读取VECIN直通的Cube计算结果
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::MatmulCompute()
{
    reluOutLocal = reluOutQueue_.AllocTensor<cType>();
    matmulObj.template GetTensorC<true>(reluOutLocal, false, true);
}

// LeakyReLU激活计算：原地逐元素运算
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::LeakyReluCompute()
{
    LeakyRelu(reluOutLocal, reluOutLocal, (cType)alpha, tiling.baseM * tiling.baseN);
    reluOutQueue_.EnQue(reluOutLocal);
}

// 数据搬出：片上结果异步拷贝至全局内存
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::CopyOut(uint32_t count)
{
    reluOutQueue_.DeQue<cType>();
    // 计算单核内迭代分块偏移
    const uint32_t roundM = tiling.singleCoreM / tiling.baseM;
    const uint32_t roundN = tiling.singleCoreN / tiling.baseN;
    uint32_t startOffset = (count % roundM * tiling.baseM * tiling.N + count / roundM * tiling.baseN);
    // 配置二维数据拷贝参数
    DataCopyParams copyParam = {(uint16_t)tiling.baseM, (uint16_t)(tiling.baseN * sizeof(cType) / DEFAULT_C0_SIZE), 0,
                                (uint16_t)((tiling.N - tiling.baseN) * sizeof(cType) / DEFAULT_C0_SIZE)};
    DataCopy(cGlobal[startOffset], reluOutLocal, copyParam);
    reluOutQueue_.FreeTensor(reluOutLocal);
}

// 多核偏移计算：匹配矩阵分块并行规则
template <typename aType, typename bType, typename cType, typename biasType>
__aicore__ inline void
MatmulLeakyKernel<aType, bType, cType, biasType>::CalcOffset(int32_t blockIdx, const TCubeTiling &tiling,
                                                             int32_t &offsetA, int32_t &offsetB, int32_t &offsetC,
                                                             int32_t &offsetBias)
{
    auto mSingleBlocks = Ceiling(tiling.M, tiling.singleCoreM);
    auto mCoreIndx = blockIdx % mSingleBlocks;
    auto nCoreIndx = blockIdx / mSingleBlocks;
    // 各张量全局内存偏移计算
    offsetA = mCoreIndx * tiling.Ka * tiling.singleCoreM;
    offsetB = nCoreIndx * tiling.singleCoreN;
    offsetC = mCoreIndx * tiling.N * tiling.singleCoreM + nCoreIndx * tiling.singleCoreN;
    offsetBias = nCoreIndx * tiling.singleCoreN;
}

// 算子内核入口函数
extern "C" __global__ __aicore__ void matmul_leakyrelu_custom(
    GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace, GM_ADDR tiling
) {
    REGISTER_TILING_DEFAULT(MatmulLeakyreluCustomTilingData);
    GET_TILING_DATA(tilingData, tiling);
    MatmulLeakyKernel<half, half, float, float> matmulLeakyKernel;
    AscendC::TPipe pipe;
    // 注册矩阵乘对象至硬件流水线
    REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulLeakyKernel.matmulObj, &tilingData.cubeTilingData);
    // 内核初始化与执行
    matmulLeakyKernel.Init(a, b, bias, c, workspace, tilingData.cubeTilingData, tilingData.alpha, &pipe);
    matmulLeakyKernel.Process(&pipe);
}

5 算子测试验证

基于ACL框架搭建端到端测试程序，完成算子初始化、内存分配、数据灌入、算子执行、结果回读与精度校验，验证融合算子功能的正确性与稳定性。

#include <algorithm>
#include <cstdint>
#include <cstdio>
#include <vector>
#include "aclnn/aclnn_base.h"
#include "aclnn/acl_meta.h"
#include "acl/acl_rt.h"
#include "aclnn_matmul_leakyrelu_custom.h"

#define CHECK_ACL(expr)                                                                                 \
    do {                                                                                                \
        auto __ret = (expr);                                                                            \
        int32_t __code = static_cast<int32_t>(__ret);                                                   \
        if (__code != 0) {                                                                              \
            fprintf(stderr, "[ERROR] %s failed at %s:%d, ret=%d\n", #expr, __FILE__, __LINE__, __code); \
        }                                                                                               \
    } while (0)

int32_t main(int32_t argc, char** argv)
{
    const int32_t deviceId = 0;
    aclrtStream stream = nullptr;
    // 初始化ACL框架与设备
    CHECK_ACL(aclnnInit(nullptr));
    CHECK_ACL(aclrtSetDevice(deviceId));
    CHECK_ACL(aclrtCreateStream(&stream));

    // 定义算子张量维度
    const std::vector<int64_t> shape_a = {1024, 256};
    const std::vector<int64_t> shape_b = {256, 640};
    const std::vector<int64_t> shape_bias = {640};
    const std::vector<int64_t> shape_output = {1024, 640};

    // 计算张量内存大小
    const int64_t elementCount_a = shape_a[0] * shape_a[1];
    const int64_t elementCount_b = shape_b[0] * shape_b[1];
    const int64_t elementCount_bias = shape_bias[0];
    const int64_t elementCount_output = shape_output[0] * shape_output[1];
    
    const size_t bufferSize_a = elementCount_a * sizeof(aclFloat16);
    const size_t bufferSize_b = elementCount_b * sizeof(aclFloat16);
    const size_t bufferSize_bias = elementCount_bias * sizeof(float);
    const size_t bufferSize_output = elementCount_output * sizeof(float);

    // 分配设备显存并创建张量
    void* inputADeviceMem = nullptr;
    CHECK_ACL(aclrtMalloc(&inputADeviceMem, bufferSize_a, ACL_MEM_MALLOC_HUGE_FIRST));
    aclTensor* inputA = aclCreateTensor(shape_a.data(), shape_a.size(), ACL_FLOAT16, nullptr, 0, ACL_FORMAT_ND,
                                        shape_a.data(), shape_a.size(), inputADeviceMem);

    void* inputBDeviceMem = nullptr;
    CHECK_ACL(aclrtMalloc(&inputBDeviceMem, bufferSize_b, ACL_MEM_MALLOC_HUGE_FIRST));
    aclTensor* inputB = aclCreateTensor(shape_b.data(), shape_b.size(), ACL_FLOAT16, nullptr, 0, ACL_FORMAT_ND,
                                        shape_b.data(), shape_b.size(), inputBDeviceMem);

    void* inputBiasDeviceMem = nullptr;
    CHECK_ACL(aclrtMalloc(&inputBiasDeviceMem, bufferSize_bias, ACL_MEM_MALLOC_HUGE_FIRST));
    aclTensor* inputBias = aclCreateTensor(shape_bias.data(), shape_bias.size(), ACL_FLOAT, nullptr, 0, ACL_FORMAT_ND,
                                           shape_bias.data(), shape_bias.size(), inputBiasDeviceMem);

    void* outputDeviceMem = nullptr;
    CHECK_ACL(aclrtMalloc(&outputDeviceMem, bufferSize_output, ACL_MEM_MALLOC_HUGE_FIRST));
    aclTensor* output = aclCreateTensor(shape_output.data(), shape_output.size(), ACL_FLOAT, nullptr, 0, ACL_FORMAT_ND,
                                        shape_output.data(), shape_output.size(), outputDeviceMem);

    // 初始化测试数据与真值
    std::vector<aclFloat16> inputAHostData(elementCount_a, aclFloatToFloat16(1.0));
    std::vector<aclFloat16> inputBHostData(elementCount_b, aclFloatToFloat16(2.0));
    std::vector<float> inputBiasHostData(elementCount_bias, float(0.5));
    std::vector<float> outputHostData(elementCount_output, float(0.0));
    std::vector<float> goldenData(elementCount_output, float(512.5));

    // 主机数据拷贝至设备
    CHECK_ACL(aclrtMemcpy(inputADeviceMem, bufferSize_a, inputAHostData.data(),
                          bufferSize_a, ACL_MEMCPY_HOST_TO_DEVICE));
    CHECK_ACL(aclrtMemcpy(inputBDeviceMem, bufferSize_b, inputBHostData.data(),
                          bufferSize_b, ACL_MEMCPY_HOST_TO_DEVICE));
    CHECK_ACL(aclrtMemcpy(inputBiasDeviceMem, bufferSize_bias, inputBiasHostData.data(),
                          bufferSize_bias, ACL_MEMCPY_HOST_TO_DEVICE));

    // 获取工作空间大小与执行器
    uint64_t workspaceSize = 0;
    aclOpExecutor* executor = nullptr;
    CHECK_ACL(aclnnMatmulLeakyreluCustomGetWorkspaceSize(inputA, inputB, inputBias, output, &workspaceSize, &executor));

    // 分配工作空间
    void* workspaceDeviceMem = nullptr;
    if (workspaceSize > 0) {
        CHECK_ACL(aclrtMalloc(&workspaceDeviceMem, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST));
    }

    // 执行自定义融合算子
    CHECK_ACL(aclnnMatmulLeakyreluCustom(workspaceDeviceMem, workspaceSize, executor, stream));
    CHECK_ACL(aclrtSynchronizeStream(stream));

    // 结果回读至主机
    CHECK_ACL(aclrtMemcpy(outputHostData.data(), bufferSize_output, outputDeviceMem,
                          bufferSize_output, ACL_MEMCPY_DEVICE_TO_HOST));

    // 结果校验与输出
    printf("算子执行结果预览：\n");
    const int64_t previewCount = std::min<int64_t>(elementCount_output, 10);
    for (int64_t i = 0; i < previewCount; i++) { 
        printf("%.1f ", outputHostData[i]); 
    }
    printf("\n测试结果：%s\n", std::equal(outputHostData.begin(), outputHostData.end(), goldenData.begin()) ? "PASS" : "FAILED");

    // 资源释放
    aclDestroyTensor(inputA);
    aclDestroyTensor(inputB);
    aclDestroyTensor(inputBias);
    aclDestroyTensor(output);
    
    CHECK_ACL(aclrtFree(inputADeviceMem));
    CHECK_ACL(aclrtFree(inputBDeviceMem));
    CHECK_ACL(aclrtFree(inputBiasDeviceMem));
    CHECK_ACL(aclrtFree(outputDeviceMem));
    
    if (workspaceSize > 0) {
        CHECK_ACL(aclrtFree(workspaceDeviceMem));
    }
    
    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclnnFinalize());
    
    return 0;
}

6 核心框架体系与理论总结

6.1 核心上下文类体系

CANN框架通过三类核心上下文类实现算子编译、分片与调度全流程管控，各司其职、分层解耦：

• TilingContext：分片调度上下文，运行于编译阶段，负责配置并行核数、申请工作空间、下发自定义分片参数，是硬件资源调度的核心入口。
• InferShapeContext：形状推导上下文，负责读取输入张量维度、修改输出张量形状，确定算子前后的张量尺寸映射关系。
• InferDataTypeContext：类型推导上下文，负责绑定算子输入输出的数据类型，保障计算精度一致性。

6.2 命名空间分层规范

• ge命名空间：顶层通用算子定义层，包含算子基类、数据类型枚举、格式枚举、执行状态码等基础定义。
• gert命名空间：编译上下文子模块，仅包含三类上下文工具类与张量工具类，服务于算子编译流程。
• matmul_tiling命名空间：矩阵乘专属分片工具层，提供多核分片、格式配置、通路配置等矩阵运算专属能力。

6.3 融合算子核心优化逻辑

VV融合通过合并连续Vector运算，消除中间数据搬移开销；CV融合依托昇腾A2分离架构，实现Cube与Vector单元的流水线并行，通过VECIN直通机制规避片上缓存数据落盘，最大化缩短计算链路时延。两类融合算子均通过Host侧参数配置+Kernel侧硬件流水线执行的分层设计，实现功能正确性与高性能的统一，可广泛应用于深度学习推理与训练的矩阵运算场景。