大家好，今天继续Ascend C 古法编程学习之旅！

Ascend C 自定义算子开发的核心，就是吃透「Host端调度 + Kernel端核内计算」的完整链路，掌握数据分片、内存搬运、核内计算、结果回写的标准化流程。上期搭建了基础开发环境，本期我将从最简单的 Hello 核函数入手，完整落地逐元素加法自定义算子 AddCustom，同时梳理通用开发模板，并预留减法 SubCustom 实战作业，后续会继续补全除法 DivCustom，帮大家打通基础二元运算算子的开发逻辑。

整篇文章我会跳出单纯代码堆砌，带着大家拆解每一段代码的设计思路、执行流程，搞懂「为什么这么写」，而非只记语法，新手也能轻松看懂！

一、开篇入门：最简 Ascend C 核函数 HelloWorld

正式开发算子前，我们先通过极简案例，熟悉 Ascend C 的核函数启动、设备初始化、流调度、资源释放整套基础流程，这是所有自定义算子的通用底座。

这段代码实现了多 Block 打印日志，直观展示 Ascend 芯片多核心并行特性，后续所有算子都会沿用这套Host端基础调度逻辑：

#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;

// 核函数：AICore 设备端执行逻辑
__global__ __aicore__ void hi_ascend(){
    // 打印当前块索引与总块数，验证多Block并行
    printf("Block[%lu/%lu]: Hi Ascend\n",GetBlockIdx(),GetBlockNum());
}

// Host端：主机侧调度入口
int32_t main(int argc ,char const *argv[]){
    // 1. 初始化ACL框架
    aclInit(nullptr);
    int32_t deviceId=0;
    aclrtSetDevice(deviceId);
    aclrtStream stream=nullptr;
    aclrtCreateStream(&stream);

    // 2. 启动4个Block并行执行核函数
    constexpr uint32_t blockDim=4;
    hi_ascend<<<blockDim,nullptr,stream>>>();

    // 3. 同步等待任务执行完成，释放资源
    aclrtSynchronizeStream(stream);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return 0;
}

核心逻辑总结：所有 Ascend C 算子都遵循「初始化设备 → 创建任务流 → 启动核函数 → 同步等待 → 释放资源」的固定流程，后续加减算子只是在这个底座上，补充了内存拷贝、数据分片、核内计算的核心逻辑。

二、核心实战：逐元素加法自定义算子 AddCustom

逐元素加法是深度学习中最基础的二元运算算子，本次我们手写完整的端到端自定义加法算子，支持 shape 为 (8, 2048) 的 float 类型张量，输入输出形状一致。通过这个案例，彻底掌握 Ascend C 算子的分层开发思想：Host 端负责调度与内存管理，Kernel 端负责核内数据搬运与计算。

2.1 前置配置与通用头文件

首先引入开发必备头文件，定义队列深度、缓存数量等全局常量，这是算子内存分片的核心配置，决定了核内数据并行搬运的能力：

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
// Host侧ACL调度接口
#include "acl/acl.h" 
// Kernel侧核函数计算接口
#include "kernel_operator.h" 
using namespace AscendC;
using namespace std;

// 核内缓存块数量、队列深度，适配流水线并行
constexpr uint32_t BUFFER_NUM=2;
constexpr uint32_t QUEUE_DEPTH=2;

2.2 分片参数结构体定义

通过结构体统一管理数据总分片参数，实现 Host 与 Kernel 端的参数传递，方便后续灵活修改分片策略：

struct AddCustomTilingData{
    uint32_t totalLength;  // 数据总长度
    uint32_t tileNum;      // 单Block内部分块数量
};

2.3 核内计算类封装（核心架构）

我们用面向对象的思想封装 KernelAdd 类，将内存初始化、数据拷贝、计算、结果回写四大核心逻辑解耦，这是 Ascend C 自定义算子的标准开发范式：

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);
    // 核心执行流程：拷贝-计算-回写
    __aicore__ inline void Process();
private:
    // 数据搬运、计算子逻辑
    __aicore__ inline void CopyIn(int32_t progress);
    __aicore__ inline void Compute(int32_t progress);
    __aicore__ inline void CopyOut(int32_t progress);

private:
    Tpipe pipe;// 流水线内存管理对象
    // 输入输出队列：管理核内局部内存张量
    TQue<TPosition::VECIN,QUEUE_DEPTH>  inQueueX,inQueueY;
    TQue<TPosition::VECOUT,QUEUE_DEPTH> outQueueZ;
    // 全局内存张量：绑定设备端全局内存
    AscendC::GlobalTensor<float> xGm,yGm,zGm;
    // 分片参数
    uint32_t blockLength;  // 单个AICore处理的数据长度
    uint32_t tileNum;      // 核内分块数
    uint32_t tileLength;   // 单块数据长度
};

2.4 初始化 Init 函数：内存资源绑定

Init 函数的核心作用是数据分片计算 + 全局/局部内存绑定初始化。根据总数据量和核心数，拆分每个 AICore 的处理任务，同时初始化流水线队列内存，为后续数据搬运做准备：

__aicore__ inline void KernelAdd::Init(GM_ADDR x,GM_ADDR y,GM_ADDR z,uint32_t totalLength,uint32_t tileNum){
    // 均分数据：每个Block处理的数据长度
    this->blockLength=totalLength/AscendC::GetBlockNum();
    this->tileNum=tileNum;
    // 细分核内分块长度，适配缓存大小
    this->tileLength=this->blockLength/tileNum/BUFFER_NUM;

    // 绑定当前Block对应的全局内存区间，避免数据重叠
    xGm.SetGlobalBuffer((__gm__ float *)x + this->blockLength*GetBlockIdx(),this->blockLength);
    yGm.SetGlobalBuffer((__gm__ float *)y +this->blockLength*GetBlockIdx(),this->blockLength);
    zGm.SetGlobalBuffer((__gm__ float *)z+this->blockLength*GetBlockIdx(),this->blockLength);

    // 初始化流水线队列内存（按字节分配）
    pipe.InitBuffer(inQueueX,BUFFER_NUM,this->tileLength*sizeof(float));
    pipe.InitBuffer(inQueueY,BUFFER_NUM,this->tileLength*sizeof(float));
    pipe.InitBuffer(outQueueZ,BUFFER_NUM,this->tileLength*sizeof(float));
}

2.5 核心执行流程 Process

采用流水线循环执行，遍历所有分块数据，依次执行「数据迁入-核内计算-结果迁出」，保证数据处理的连续性和高效性：

__aicore__ inline void KernelAdd::Process(){
    // 总循环次数 = 分块数 * 缓存块数
    int32_t loopCount=this->tileNum*BUFFER_NUM;
    for(int32_t i=0;i<loopCount;i++){
        CopyIn(i);   // 全局内存 -> 核内局部内存
        Compute(i);  // 核内逐元素加法计算
        CopyOut(i);  // 核内局部内存 -> 全局内存
    }
}

2.6 数据搬运与计算逻辑

这是算子的计算核心，三段逻辑各司其职：CopyIn 负责将设备全局内存数据搬运到 AICore 局部内存；Compute 调用 Ascend C 内置加法算子完成计算；CopyOut 将计算结果写回全局内存，同时释放无用内存，避免显存泄漏。

（已修正原版代码拼写错误，保证可直接编译运行）

// 数据迁入：全局内存转局部内存
__aicore__ inline void KernelAdd::CopyIn(int32_t progress){
    LocalTensor<float> xLocal=inQueueX.AllocTensor<float>();
    LocalTensor<float> yLocal=inQueueY.AllocTensor<float>();
    DataCopy(xLocal,xGm[progress*this->tileLength],this->tileLength);
    DataCopy(yLocal,yGm[progress*this->tileLength],this->tileLength);
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}

// 核内加法计算
__aicore__ inline void KernelAdd::Compute(int32_t progress){
    LocalTensor<float> xLocal=inQueueX.DeQue<float>();
    LocalTensor<float> yLocal=inQueueY.DeQue<float>();
    LocalTensor<float> zLocal=outQueueZ.AllocTensor<float>();
    // 逐元素加法核心接口
    Add(zLocal,xLocal,yLocal,this->tileLength);
    outQueueZ.EnQue<float>(zLocal);
    // 释放内存，优化显存占用
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

// 结果迁出：局部内存写回全局内存
__aicore__ inline void KernelAdd::CopyOut(int32_t progress){
    LocalTensor<float> zLocal=outQueueZ.DeQue<float>();
    DataCopy(zGm[progress*tileLength],zLocal,tileLength);
    outQueueZ.FreeTensor(zLocal);
}

2.7 核函数入口

全局核函数入口，指定任务类型、实例化计算类、启动初始化与计算流程，是 Host 端与 Kernel 端的衔接入口：

__global__ __aicore__ void add_custom(GM_ADDR x,GM_ADDR y,GM_ADDR z,AddCustomTilingData tiling){
    // 声明纯AI计算任务
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    KernelAdd op;
    op.Init(x,y,z,tiling.totalLength,tiling.tileNum);
    op.Process();
}

2.8 Host 端调度与内存管理

Host 端核心职责：初始化设备、申请主机/设备内存、数据拷贝、启动核函数、结果回读、资源释放，是整个算子的调度中枢：

vector<float> kernel_add(vector<float> &x,vector<float> &y){
    constexpr uint32_t blockDim=8;
    uint32_t totalLength=x.size();
    size_t totalByteSize=totalLength*sizeof(float);
    int32_t deviceId=0;
    aclrtStream stream=nullptr;
    // 初始化分片参数
    AddCustomTilingData tiling={totalLength,8};

    // 主机、设备内存指针定义
    uint8_t *xHost=reinterpret_cast<uint8_t *>(x.data());
    uint8_t *yHost=reinterpret_cast<uint8_t *>(y.data());
    uint8_t *zHost=nullptr,*xDevice=nullptr,*yDevice=nullptr,*zDevice=nullptr;

    // 设备初始化
    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreateStream(&stream);

    // 内存申请
    aclrtMallocHost((void**)(&zHost),totalByteSize);
    aclrtMalloc((void **)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);

    // 数据：主机 -> 设备
    aclrtMemcpy(xDevice, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(yDevice, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

    // 启动自定义加法核函数
    add_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
    aclrtSynchronizeStream(stream);

    // 结果：设备 -> 主机
    aclrtMemcpy(zHost, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float *)zHost, (float *)(zHost + totalByteSize));

    // 资源统一释放
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return z;
}

2.9 结果验证与主函数

编写通用验证函数，对比算子输出结果与标准真值，快速校验算子正确性，主函数构造测试数据、调用算子、完成验证：

// 结果验证函数
uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
{
    auto printTensor = [](std::vector<float> &tensor, const char *name) {
        constexpr size_t maxPrintSize = 20;
        std::cout << name << ": ";
        std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
            std::ostream_iterator<float>(std::cout, " "));
        if (tensor.size() > maxPrintSize) std::cout << "...";
        std::cout << std::endl;
    };
    printTensor(output, "Output");
    printTensor(golden, "Golden");
    if (std::equal(golden.begin(), golden.end(), output.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
}

// 主测试函数
int32_t main(int32_t argc, char *argv[])
{
    // 适配固定shape (8,2048)
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    // 构造测试张量
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);
    // 调用自定义加法算子
    std::vector<float> output = kernel_add(x, y);
    // 标准真值
    std::vector<float> golden(totalLength, valueX + valueY);
    return VerifyResult(output, golden);
}

三、课后实战：复刻减法算子 SubCustom（填空拓展）

掌握加法算子后，减法算子是完全同架构复刻的拓展练习，整体流程、内存管理、分片逻辑与加法一致，仅需替换核心计算接口 Add -> Sub。

我保留了完整代码框架，修正了原版所有语法错误（拼写错误、参数缺失、指针错误等），大家可基于框架补全代码，亲手巩固整套算子开发流程。

实战要求：

• 数据类型：float
• 数据Shape：(8, 2048)，输入输出形状一致
• 数据布局：ND 标准布局

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;
using namespace std;

// 全局流水线配置，与加法算子保持一致
constexpr uint32_t BUFFER_NUM = 2;
constexpr uint32_t QUEUE_DEPTH = 2;
constexpr uint32_t TILE_NUM = 8;
constexpr uint32_t BLOCK_NUM = 8;

// 减法算子分片参数结构体
struct SubCustomTilingData
{
    uint32_t totalLength;
    uint32_t tileNum;
};

// 逐元素减法核计算类
class KernelSub {
public:
    __aicore__ inline KernelSub(){}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum);
    __aicore__ inline void Process();

private:
    __aicore__ inline void CopyIn(int32_t progress);
    __aicore__ inline void Compute(int32_t progress);
    __aicore__ inline void CopyOut(int32_t progress);

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

// 初始化：数据分片 + 全局内存绑定 + 流水线队列初始化
__aicore__ inline void KernelSub::Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
{
    // 按Block均分数据，单核心处理数据长度
    this->blockLength = totalLength / AscendC::GetBlockNum();
    this->tileNum = tileNum;
    // 细分核内小分块长度，适配缓存流水线
    this->tileLength = this->blockLength / tileNum / BUFFER_NUM;

    // 绑定当前Block对应的全局内存区间，数据互不干扰
    xGm.SetGlobalBuffer((__gm__ float *)x + this->blockLength * GetBlockIdx(), this->blockLength);
    yGm.SetGlobalBuffer((__gm__ float *)y + this->blockLength * GetBlockIdx(), this->blockLength);
    zGm.SetGlobalBuffer((__gm__ float *)z + this->blockLength * GetBlockIdx(), this->blockLength);

    // 初始化输入输出队列内存空间
    pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
    pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
    pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
}

// 整体流水线执行逻辑
__aicore__ inline void KernelSub::Process()
{
    int32_t loopCount = tileNum * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++)
    {
        CopyIn(i);
        Compute(i);
        CopyOut(i);
    }
}

// 数据迁入：全局内存 -> AICore局部内存
__aicore__ inline void KernelSub::CopyIn(int32_t progress)
{
    LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
    LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
    DataCopy(xLocal, xGm[progress * tileLength], tileLength);
    DataCopy(yLocal, yGm[progress * tileLength], tileLength);
    inQueueX.EnQue(xLocal);
    inQueueY.EnQue(yLocal);
}

// 核内减法计算核心逻辑
__aicore__ inline void KernelSub::Compute(int32_t progress)
{
    LocalTensor<float> xLocal = inQueueX.DeQue<float>();
    LocalTensor<float> yLocal = inQueueY.DeQue<float>();
    LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
    // 逐元素减法核心接口 z = x - y
    Sub(zLocal, xLocal, yLocal, tileLength);
    outQueueZ.EnQue<float>(zLocal);
    // 释放无用局部内存
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

// 结果迁出：局部计算结果 -> 全局内存
__aicore__ inline void KernelSub::CopyOut(int32_t progress)
{
    LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
    DataCopy(zGm[progress * tileLength], zLocal, tileLength);
    outQueueZ.FreeTensor(zLocal);
}

// 减法核函数全局入口
__global__ __aicore__ void sub_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, SubCustomTilingData tiling)
{
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    KernelSub op;
    op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
    op.Process();
}

// Host端调度封装：内存管理 + 核函数启动 + 数据交互
std::vector<float> kernel_sub(std::vector<float> &x, std::vector<float> &y)
{
    constexpr uint32_t blockDim = BLOCK_NUM;
    uint32_t totalLength = x.size();
    size_t totalByteSize = totalLength * sizeof(float);
    int32_t deviceId = 0;
    aclrtStream stream = nullptr;
    SubCustomTilingData tiling = {totalLength, TILE_NUM};

    uint8_t *xHost = reinterpret_cast<uint8_t*>(x.data());
    uint8_t *yHost = reinterpret_cast<uint8_t*>(y.data());
    uint8_t *zHost = nullptr;
    uint8_t *xDevice = nullptr;
    uint8_t *yDevice = nullptr;
    uint8_t *zDevice = nullptr;

    // ACL设备初始化
    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreateStream(&stream);

    // 申请主机、设备显存
    aclrtMallocHost((void**)(&zHost), totalByteSize);
    aclrtMalloc((void**)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void**)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void**)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);

    // 数据从主机拷贝至设备
    aclrtMemcpy(xDevice, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(yDevice, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

    // 启动自定义减法核函数
    sub_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
    aclrtSynchronizeStream(stream);

    // 计算结果回读到主机内存
    aclrtMemcpy(zHost, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float*)zHost, (float*)(zHost + totalByteSize));

    // 统一释放所有资源，防止内存泄漏
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();

    return z;
}

// 通用结果验证函数
uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
{
    auto printTensor = [](std::vector<float> &tensor, const char *name) {
        constexpr size_t maxPrintSize = 20;
        std::cout << name << ": ";
        std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
            std::ostream_iterator<float>(std::cout, " "));
        if (tensor.size() > maxPrintSize) std::cout << "...";
        std::cout << std::endl;
    };
    printTensor(output, "Output");
    printTensor(golden, "Golden");
    if (std::equal(golden.begin(), golden.end(), output.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
}

// 主测试入口
int32_t main(int32_t argc, char *argv[])
{
    // 固定测试Shape (8, 2048)
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);

    // 调用自定义减法算子
    std::vector<float> output = kernel_sub(x, y);
    // 标准真值：x - y
    std::vector<float> golden(totalLength, valueX - valueY);

    return VerifyResult(output, golden);
}

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;
constexpr uint32_t BUFFER_NUM = 2; // tensor num for each queue
constexpr uint32_t QUEUE_DEPTH = 2;
constexpr uint32_t TILE_NUM=8;
constexpr uint32_t BLOCK_NUM=8;

struct SubCustomTilingData
{
    uint32_t totalLength;
    uint32_t tileNum;
};

class KernelSub {
public:
    __aicore__ inline KernelSub(){}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
    {
        // 请补充……
        
    }
    __aicore__ inline void Process()
    {
        // 请补充……
        
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        // 请补充……
        
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        // 请补充……
        
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        // 请补充……
        
    }

private:
    // 请补充……
    Tpipe pipe;
    TQue<AscendC::TPosition::VECIN,QUEUE_DEPTH> inQueueX,inQueueY;
    TQue<TPosition::VECOUT,QUEUE_DEPTH> outQueueZ;
    GlobalTensor<float> xGm,yGm,zGm;
    uint32_t blockLength,tileNum,tileLength;//tiling结构体 totalLength，tileNum
};

__global__ __aicore__ void sub_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, SubCustomTilingData tiling)
{
    // 请补充……
    
}

std::vector<float> kernel_sub(std::vector<float> &x, std::vector<float> &y)
{
    // 请补充……
    
}

uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
{
    auto printTensor = [](std::vector<float> &tensor, const char *name) {
        constexpr size_t maxPrintSize = 20;
        std::cout << name << ": ";
        std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
            std::ostream_iterator<float>(std::cout, " "));
        if (tensor.size() > maxPrintSize) {
            std::cout << "...";
        }
        std::cout << std::endl;
    };
    printTensor(output, "Output");
    printTensor(golden, "Golden");
    if (std::equal(golden.begin(), golden.end(), output.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
    return 0;
}

int32_t main(int32_t argc, char *argv[])
{
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);

    // 请补充……
    std::vector<float> output=kernel_sub(x,y);

    std::vector<float> golden(totalLength, valueX - valueY);
    return VerifyResult(output, golden);
}

完整参考实现

以下为对齐加法算子架构、修复全部BUG后的SubCustom 完整源码，可对照填空练习自查、运行验证：

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;
using namespace std;

constexpr uint32_t BUFFER_NUM = 2;
constexpr uint32_t QUEUE_DEPTH = 2;
constexpr uint32_t TILE_NUM = 8;
constexpr uint32_t BLOCK_NUM = 8;

struct SubCustomTilingData
{
    uint32_t totalLength;
    uint32_t tileNum;
};

class KernelSub {
public:
    __aicore__ inline KernelSub(){}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
    {
        this->blockLength = totalLength / AscendC::GetBlockNum();
        this->tileNum = tileNum;
        this->tileLength = this->blockLength / tileNum / BUFFER_NUM;

        xGm.SetGlobalBuffer((__gm__ float *)x + this->blockLength * GetBlockIdx(), this->blockLength);
        yGm.SetGlobalBuffer((__gm__ float *)y + this->blockLength * GetBlockIdx(), this->blockLength);
        zGm.SetGlobalBuffer((__gm__ float *)z + this->blockLength * GetBlockIdx(), this->blockLength);

        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
        pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
    }

    __aicore__ inline void Process()
    {
        int32_t loopCount = 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)
    {
        LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
        LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
        DataCopy(xLocal, xGm[progress * tileLength], tileLength);
        DataCopy(yLocal, yGm[progress * tileLength], tileLength);
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }

    __aicore__ inline void Compute(int32_t progress)
    {
        LocalTensor<float> xLocal = inQueueX.DeQue<float>();
        LocalTensor<float> yLocal = inQueueY.DeQue<float>();
        LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();

        Sub(zLocal, xLocal, yLocal, tileLength);

        outQueueZ.EnQue<float>(zLocal);
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }

    __aicore__ inline void CopyOut(int32_t progress)
    {
        LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
        DataCopy(zGm[progress * tileLength], zLocal, tileLength);
        outQueueZ.FreeTensor(zLocal);
    }

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

__global__ __aicore__ void sub_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, SubCustomTilingData tiling)
{
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    KernelSub op;
    op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
    op.Process();
}

std::vector<float> kernel_sub(std::vector<float> &x, std::vector<float> &y)
{
    constexpr uint32_t blockDim = BLOCK_NUM;
    uint32_t totalLength = x.size();
    size_t totalByteSize = totalLength * sizeof(float);
    int32_t deviceId = 0;
    aclrtStream stream = nullptr;
    SubCustomTilingData tiling = {totalLength, TILE_NUM};

    uint8_t *xHost = reinterpret_cast<uint8_t*>(x.data());
    uint8_t *yHost = reinterpret_cast<uint8_t*>(y.data());
    uint8_t *zHost = nullptr;
    uint8_t *xDevice = nullptr;
    uint8_t *yDevice = nullptr;
    uint8_t *zDevice = nullptr;

    aclInit(nullptr);
    aclrtSetDevice(deviceId);
    aclrtCreateStream(&stream);

    aclrtMallocHost((void**)(&zHost), totalByteSize);
    aclrtMalloc((void**)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void**)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void**)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);

    aclrtMemcpy(xDevice, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(yDevice, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);

    sub_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
    aclrtSynchronizeStream(stream);

    aclrtMemcpy(zHost, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float*)zHost, (float*)(zHost + totalByteSize));

    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();

    return z;
}

uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
{
    auto printTensor = [](std::vector<float> &tensor, const char *name) {
        constexpr size_t maxPrintSize = 20;
        std::cout << name << ": ";
        std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
            std::ostream_iterator<float>(std::cout, " "));
        if (tensor.size() > maxPrintSize) {
            std::cout << "...";
        }
        std::cout << std::endl;
    };
    printTensor(output, "Output");
    printTensor(golden, "Golden");
    if (std::equal(golden.begin(), golden.end(), output.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
    return 0;
}

int32_t main(int32_t argc, char *argv[])
{
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);

    std::vector<float> output = kernel_sub(x, y);
    std::vector<float> golden(totalLength, valueX - valueY);

    return VerifyResult(output, golden);
}

四、学习总结与下期预告

本期我们彻底吃透了Ascend C 自定义二元算子的通用开发模板：

✅ 掌握 Host 与 Kernel 端的分层协作逻辑

✅ 熟悉「分片初始化-数据迁入-核内计算-结果回写」流水线

✅ 独立落地完整的 AddCustom 加法算子，修复各类实操报错

✅ 完成 SubCustom 减法算子填空练习 + 完整源码落地

下期我会补全 DivCustom 除法自定义算子，完善加减除基础算子体系，同时对比四则运算算子的开发异同，帮大家固化开发思维，后续可以快速适配各类逐元素运算算子！

（注：部分内容可能由 AI 生成）