前言

本期我们梳理了算子开发的完整开发链路，吃透了「Host 主机调度 + AICore Kernel 核内流水线」标准化开发框架，进阶学习双输入二元逐元素算子，先从最简 HelloWorld 熟悉设备调度基础流程，再完整落地逐元素加法算子 AddCustom，拆解整套通用开发模板；同时配套 SubCustom 填空实战作业，并附上可直接编译运行的完整参考源码。 整篇内容使用开发教程的写作逻辑，不单纯堆砌代码，每段配套底层原理注释、高频疑问解答，拆解分片、内存拷贝、流水线双缓冲、Host/Device 内存交互核心逻辑，新手也能看懂底层设计思路，吃透 “为什么这么写”，而不是死记硬背代码模板。

一、开发前置认知

所有 Ascend C 自定义算子遵循统一分层架构：

  1. Kernel 设备侧：运行在昇腾 AICore，负责数据分片搬运、本地缓存流水线、内置算子计算，分为 Init/Process/CopyIn/Compute/CopyOut 标准五段式结构；
  2. Host 主机侧：运行在 CPU，负责 ACL 设备初始化、Host/Device 内存分配、双向数据拷贝、核函数异步启动、流同步、资源释放；
  3. 统一流水线设计：采用双缓冲 BUFFER_NUM=2 实现计算与数据拷贝并行，隐藏 IO 耗时，提升硬件利用率；
  4. 统一验证体系：CPU 标准真值函数 + 批量打印校验函数，快速验证算子精度。

本次算子统一规范：

  • 数据类型：输入输出均为 float
  • 张量 Shape：固定 (8, 2048)，一维展开总长度 8*2048
  • 数据布局：ND 标准稠密布局
  • 并行策略：启动 8 个 Block 均分全部数据，单核内部再细分多块 tile 流水线处理

二、开篇入门：HelloWorld 极简核函数（通用调度底座）

正式开发算子前，先通过极简示例掌握昇腾程序固定执行流程，所有加减乘除、激活算子都会复用这套 Host 基础逻辑：初始化 ACL、创建设备流、启动核函数、同步等待、释放全部资源。 这段代码多 Block 打印日志，直观理解 AICore 多核心并行特性。

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

// Kernel全局核函数：运行在AICore设备端
__global__ __aicore__ void hi_ascend(){
    // 打印当前块索引、总块数量，验证多核心并行拆分效果
    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;
}

基础流程总结

全系列算子通用执行链路： 初始化设备 → 创建异步任务流 <<<>>> 启动核函数 → stream同步等待计算结束 → 销毁流/重置设备/释放ACL 后续 Add、Sub 算子仅在此基础上增加张量内存分配、分片参数传递、GM 与 LocalTensor 数据搬运、核内数学计算模块。

三、核心实战：逐元素加法算子 AddCustom 端到端完整实现

逐元素加法是深度学习最基础二元算子，输入两个同 shape 张量 x、y，输出 z=x+y。我们完整实现 Kernel 设备侧 + Host 主机侧代码，拆解每一块的设计目的。

3.1 全局常量与通用头文件

头文件区分 Host 侧 ACL 接口、Kernel 侧算子接口；全局常量控制流水线缓存配置，和 Sigmoid 算子保持统一规范。

#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
// Host主机ACL运行时接口
#include "acl/acl.h" 
// AICore核内计算底层API
#include "kernel_operator.h" 
using namespace AscendC;
using namespace std;

// 流水线核心配置：双缓冲并行
constexpr uint32_t BUFFER_NUM=2;    // 队列缓存张量数量，一边拷贝一边计算
constexpr uint32_t QUEUE_DEPTH=2;   // TQue队列深度，控制异步读写容量

3.2 Tiling 分片参数结构体

Host 通过结构体把全局数据长度、单核分块数量传递到 Kernel，统一管理分片逻辑，后续修改并行粒度只需要调整入参，不用改动核内代码。

struct AddCustomTilingData{
    uint32_t totalLength;  // 全部输入张量一维总长度
    uint32_t tileNum;      // 单个AICore内部细分小分块数量
};

3.3 KernelAdd 核计算封装类（标准五段式架构）

面向对象封装所有核内逻辑，解耦内存初始化、数据迁入、计算、结果写出，是 Ascend C 自定义算子标准范式，和 SigmoidKernel 类结构完全对齐。

class KernelAdd{
public:
    // 空构造
    __aicore__ inline KernelAdd(){}
    /// @brief Init初始化：分片长度计算 + GlobalTensor全局内存绑定 + 流水线队列内存分配
    /// @param x/y/z 输入输出全局内存地址
    /// @param totalLength 张量总长度
    /// @param tileNum 单核细分块数量
    __aicore__ inline void Init(GM_ADDR x,GM_ADDR y,GM_ADDR z,uint32_t totalLength,uint32_t tileNum);
    /// @brief Process流水线总调度入口：循环执行CopyIn -> Compute -> CopyOut
    __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;          // 输出向量队列
    GlobalTensor<float> xGm,yGm,zGm;                // 绑定设备全局GM内存张量
    uint32_t blockLength;  // 单个Block(AICore)独占处理数据长度
    uint32_t tileNum;      // 单核内部细分块数
    uint32_t tileLength;   // 单个小分块数据长度
};

3.4 Init 初始化函数详解

核心完成两件事：数据均匀分片、绑定当前 Block 专属内存区间、为输入输出队列分配 Local 本地缓存字节空间。

__aicore__ inline void KernelAdd::Init(GM_ADDR x,GM_ADDR y,GM_ADDR z,uint32_t totalLength,uint32_t tileNum){
    // 1. 均分全局数据：总长度 / 并行Block数量，每个核心只处理自己对应的区间
    this->blockLength=totalLength/AscendC::GetBlockNum();
    this->tileNum=tileNum;
    // 细分小块长度，适配AICore有限本地缓存，BUFFER_NUM双缓冲拆分
    this->tileLength=this->blockLength/tileNum/BUFFER_NUM;

    // 2. 绑定当前Block专属GM内存，GetBlockIdx获取当前核心编号，多核心数据互不重叠
    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);

    // 3. 流水线队列分配内存，入参BUFFER_NUM代表双缓冲，size按字节计算
    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));
}
高频疑问：InitBuffer 第二个参数为什么是 BUFFER_NUM？

BUFFER_NUM 代表队列内部缓存的 LocalTensor 数量，取值 2 实现双缓冲流水线：一块本地缓存正在执行计算，另一块同步从 GM 搬运下一批数据，掩盖内存拷贝延迟，最大化 AICore 算力利用率；单输入 / 双输入算子统一配置 BUFFER_NUM=2，架构通用。

3.5 Process 流水线循环调度

总循环次数 = 单核分块数 × 双缓冲数量，循环依次执行「载入数据 - 计算 - 写出结果」完整链路。

__aicore__ inline void KernelAdd::Process(){
    int32_t loopCount=this->tileNum*BUFFER_NUM;
    for(int32_t i=0;i<loopCount;i++){
        CopyIn(i);   // GM全局内存 → AICore Local本地内存
        Compute(i);  // 本地张量逐元素加法计算
        CopyOut(i);  // Local计算结果 → GM全局输出内存
    }
}

3.6 CopyIn / Compute / CopyOut 三段核心逻辑

// CopyIn：从全局内存搬运单块数据到本地输入队列
__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);
}

// Compute：取出两路本地张量，调用AscendC内置Add逐元素求和
__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);
    // 释放无用本地张量，节省L0/L1缓存
    inQueueX.FreeTensor(xLocal);
    inQueueY.FreeTensor(yLocal);
}

// CopyOut：计算完成的本地张量写回设备全局输出内存
__aicore__ inline void KernelAdd::CopyOut(int32_t progress){
    LocalTensor<float> zLocal=outQueueZ.DeQue<float>();
    DataCopy(zGm[progress*tileLength],zLocal,tileLength);
    outQueueZ.FreeTensor(zLocal);
}

3.7 全局核函数入口

Host 端 <<<>>> 调度的唯一入口，声明任务类型为纯 AI 计算，实例化 KernelAdd 对象并启动流水线。

__global__ __aicore__ void add_custom(GM_ADDR x,GM_ADDR y,GM_ADDR z,AddCustomTilingData tiling){
    KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
    KernelAdd op;
    op.Init(x,y,z,tiling.totalLength,tiling.tileNum);
    op.Process();
}

3.8 Host 侧 kernel_add 调度函数（内存交互核心）

Host 侧完整链路：初始化设备 → 分配锁页 Host 内存、设备 GM 显存 → CPU 数据拷贝至芯片显存 → 异步启动核函数 → 同步等待 → 结果拷贝回 CPU 内存 → 统一释放全部资源。 配套解释高频疑问：totalByteSize、uint8_t 强转、Host/Device 内存区分、aclrtMemcpy 参数规范。

vector<float> kernel_add(vector<float> &x,vector<float> &y){
    constexpr uint32_t blockDim=8;
    uint32_t totalLength=x.size();
    // totalByteSize：内存拷贝API仅识别字节流，总字节=元素个数×单float4字节
    size_t totalByteSize=totalLength*sizeof(float);
    int32_t deviceId=0;
    aclrtStream stream=nullptr;
    AddCustomTilingData tiling={totalLength,8};

    // uint8_t* 强转原因：ACL内存接口底层操作字节流，不感知float/int类型，统一用1字节指针管理整块内存
    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);

    // 分配CPU锁页Host内存、昇腾设备全局显存
    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 4参标准同步接口规范：dst, src, size, kind
    参数顺序口诀：目标，来源，字节长度，拷贝方向
    禁止错误写法：dst, size, src, size, kind（偏移参数与长度混淆，会内存越界崩溃）
    */
    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);

    // 设备显存计算结果拷贝回CPU主机内存
    aclrtMemcpy(zHost, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> z((float *)zHost, (float *)(zHost + totalByteSize));

    // Host与Device内存物理隔离，必须统一释放防止内存泄漏
    aclrtFree(xDevice);
    aclrtFree(yDevice);
    aclrtFree(zDevice);
    aclrtFreeHost(zHost);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    return z;
}
核心疑问统一解答
  1. totalByteSize 作用：aclrtMemcpy 底层操作字节，无 float 类型概念，必须计算整块内存总字节数作为拷贝长度；
  2. uint8_t 指针强转：统一字节流操作，适配 ACL 底层内存接口，避免类型截断；
  3. Host/Device 内存区分：Host=CPU 内存，仅 CPU 读写；Device = 昇腾芯片显存，仅 AICore 访问；两者物理隔离，必须通过 aclrtMemcpy 完成数据交换，不能直接跨端指针访问；
  4. aclrtMemcpy 参数：日常单块张量使用 4 参同步接口，顺序dst,src,size,kind；7 参异步接口才携带偏移量，不可混淆参数位置。

3.9 通用校验函数 + 测试主程序

和 Sigmoid 教程复用同一套校验逻辑，打印前 20 个数值对比算子输出与 CPU 标准真值，快速定位精度问题。

// 通用精度验证函数
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[])
{
    constexpr uint32_t totalLength = 8 * 2048;
    constexpr float valueX = 1.2f;
    constexpr float valueY = 2.3f;
    // 构造固定shape测试张量
    std::vector<float> x(totalLength, valueX);
    std::vector<float> y(totalLength, valueY);
    // 调用昇腾自定义加法算子
    std::vector<float> output = kernel_add(x, y);
    // CPU标准真值 z = x + y
    std::vector<float> golden(totalLength, valueX + valueY);
    return VerifyResult(output, golden);
}

四、课后实战：SubCustom 逐元素减法算子（填空练习框架）

掌握加法算子后，减法算子属于同架构复刻拓展，整体分片、流水线、内存交互逻辑完全不变，仅将 Compute 内部Add底层算子替换为Sub，输入输出张量数量、缓存配置、Host 调度逻辑全部复用。 下方保留填空练习框架，标注// 请补充……，适合手动填空巩固开发流程。

实战需求

  • 数据类型：float
  • Shape：(8,2048)，输入输出同 shape
  • 数据布局：ND 稠密布局
  • 计算公式：z = x - y
#include <cstdint>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iterator>
#include "acl/acl.h"
#include "kernel_operator.h"
using namespace AscendC;
using namespace std;

// 全局流水线配置，与加法、Sigmoid算子统一
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;
};

// 减法全局核函数入口
__global__ __aicore__ void sub_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, SubCustomTilingData tiling)
{
    // 请补充……
}

// Host侧减法调度封装
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);
}

SubCustom 完整可运行参考实现（填空对照源码）

下方为对齐 AddCustom 架构、修复全部语法、参数错误的完整版代码，写完填空后可直接对照自查、编译运行。

#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>();
        // 逐元素减法核心接口 z = x - y
        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);
}

五、学习总结 & 下期预告

本期核心知识点梳理

  1. 统一开发模板固化：和 Sigmoid 单输入算子对齐架构，双输入二元算子依旧遵循 Init分片 → Process流水线 → CopyIn/Compute/CopyOut 标准五段式开发范式；
  2. 流水线双缓冲底层原理：BUFFER_NUM=2 实现计算与数据拷贝并行，解决 AICore 本地缓存不足、IO 阻塞问题；
  3. Host-Device 完整内存链路：锁页内存 / 显存分配、4 参标准 aclrtMemcpy 规范、Host 与 Device 内存物理隔离核心概念；
  4. 二元算子复用逻辑：Add、Sub 算子 90% 代码完全通用，仅替换 Compute 内部底层数学 API，拓展性极强；
  5. 标准化验证体系：统一通用 VerifyResult 校验函数，CPU 生成标准真值快速校验算子精度。

下期拓展预告

当前已掌握双输入二元加法 AddCustom、减法 SubCustom，下一期将基于同一套通用模板完整实现单输入激活算子 SigmoidCustom，并逐步拓展DivCustom 逐元素除法算子，完善加减除基础二元算子体系，横向对比三类二元算子代码差异，彻底吃透 Ascend C 逐元素算子通用开发思维，后续可快速拓展 Relu、Tanh、Mul 等各类算子。

Logo
鲲鹏昇腾开发者社区是面向全社会开放的“联接全球计算开发者，聚合华为+生态”的社区，内容涵盖鲲鹏、昇腾资源，帮助开发者快速获取所需的知识、经验、软件、工具、算力，支撑开发者易学、好用、成功，成为核心开发者。

鲲鹏昇腾开发者社区是面向全社会开放的“联接全球计算开发者，聚合华为+生态”的社区，内容涵盖鲲鹏、昇腾资源，帮助开发者快速获取所需的知识、经验、软件、工具、算力，支撑开发者易学、好用、成功，成为核心开发者。

更多推荐

Ascend C 算子实战（二）｜SigmoidCustom 逐元素激活算子完整开发指南

Init分片初始化 → Process循环调度 → CopyIn/Compute/CopyOut，单输入 / 双输入算子框架完全通用，仅增减输入队列与计算 API；掌握 Ascend C 基础数学算子串联：Muls 标量乘、Exp 指数、Adds 标量加、Duplicate 填充、Div 逐元素除法；理清 Host-Device 内存交互完整链路：锁页内存申请、显存分配、双向拷贝、同步等待、资源释

avatar 鲲鹏昇腾开发者社区

【NVIDIA】Ascend Docker Runtime v26.0.1 之五 install/ — 超深度逐行分析

Ascend Docker Runtime v26.0.1 之五 install/ — 超深度逐行分析

avatar 鲲鹏昇腾开发者社区

【NPU】Ascend Docker Runtime v26.0.1 之六 mindxcheckutils/ — 超深度逐行分析

Ascend Docker Runtime v26.0.1 之六 mindxcheckutils/ — 超深度逐行分析

avatar 鲲鹏昇腾开发者社区