完整代码地址：https://download.csdn.net/download/feng8403000/92484245

概述

Add算子是深度学习中最基础的元素级算子之一，用于实现两个张量的逐元素相加。虽然Add算子看似简单，但它的实现涉及Ascend C编程的核心概念，包括内存管理、数据拷贝、向量化计算等。本文将从最基础的实现开始，逐步深入，展示如何在0_helloworld项目基础上实现一个完整的Add算子。

整体运行效果：

验证效果：

什么是Add算子

Add算子（Addition Operator）是元素级算子（Element-wise Operator）的一种，它对两个输入张量的对应位置元素进行相加运算，生成输出张量。数学表达式为：

output[i] = input1[i] + input2[i]

其中，i表示元素在张量中的索引位置。

Add算子的特点

1. 元素独立性：每个输出元素只依赖于对应位置的输入元素，元素之间没有依赖关系
2. 易于并行化：由于元素独立性，可以充分利用多核并行计算
3. 易于向量化：可以使用向量指令同时处理多个元素
4. 内存访问模式简单：顺序访问，缓存友好

Add算子的应用场景

• 残差连接（Residual Connection）：在ResNet等网络中使用
• 特征融合：将多个特征图相加融合
• 偏置添加：在卷积或全连接层后添加偏置
• 广播加法：支持不同形状张量的广播相加

基于0_helloworld实现Add算子

我们将基于0_helloworld项目来实现Add算子。0_helloworld项目提供了最基础的核函数调用框架，我们只需要修改核函数实现部分即可。

项目结构

在0_helloworld项目基础上，我们需要修改以下文件：

0_helloworld/
├── CMakeLists.txt          # 编译配置文件（基本不变）
├── hello_world.cpp         # 改为 add_custom.cpp（核函数实现）
├── main.cpp                # 修改主程序（添加数据准备和验证）
└── run.sh                  # 运行脚本（基本不变）

第一步：核函数实现（hello_world.cpp）

在Ascend C中，LocalTensor不能直接通过Alloc()方法分配，必须使用TPipe和TQue来管理内存。这是Ascend C的标准做法，可以更好地管理内存和实现流水线优化。

/**
 * @file hello_world.cpp
 * 
 * Add算子实现 - 基于0_helloworld项目修改
 */
#include "kernel_operator.h"

constexpr uint32_t TOTAL_LENGTH = 2048;

/**
 * Add算子Kernel类
 * 使用TPipe和TQue来管理LocalTensor的内存分配
 */
class KernelAdd {
public:
    __aicore__ inline KernelAdd() {}
    
    /**
     * 初始化函数
     * @param x 第一个输入张量的全局内存地址
     * @param y 第二个输入张量的全局内存地址
     * @param z 输出张量的全局内存地址
     */
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
    {
        // 1. 创建GlobalTensor对象，绑定全局内存
        xGm.SetGlobalBuffer((__gm__ half *)x, TOTAL_LENGTH);
        yGm.SetGlobalBuffer((__gm__ half *)y, TOTAL_LENGTH);
        zGm.SetGlobalBuffer((__gm__ half *)z, TOTAL_LENGTH);

        // 2. 初始化TPipe和TQue，用于管理LocalTensor的内存
        // InitBuffer会为队列分配内存空间
        pipe.InitBuffer(inQueueX, 1, TOTAL_LENGTH * sizeof(half));
        pipe.InitBuffer(inQueueY, 1, TOTAL_LENGTH * sizeof(half));
        pipe.InitBuffer(outQueueZ, 1, TOTAL_LENGTH * sizeof(half));
    }
    
    /**
     * 处理函数，执行完整的Add算子流程
     */
    __aicore__ inline void Process()
    {
        CopyIn();   // 从全局内存拷贝到本地内存
        Compute();  // 执行Add计算
        CopyOut();  // 从本地内存拷贝回全局内存
    }

private:
    /**
     * CopyIn阶段：从全局内存拷贝数据到本地内存
     */
    __aicore__ inline void CopyIn()
    {
        // 从队列中分配LocalTensor（内存由TQue管理）
        AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.AllocTensor<half>();

        // 从GlobalTensor拷贝到LocalTensor
        AscendC::DataCopy(xLocal, xGm, TOTAL_LENGTH);
        AscendC::DataCopy(yLocal, yGm, TOTAL_LENGTH);

        // 将LocalTensor放入队列（用于后续的Compute阶段）
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
    }
    
    /**
     * Compute阶段：执行Add计算
     */
    __aicore__ inline void Compute()
    {
        // 从队列中取出LocalTensor
        AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
        
        // 为输出分配LocalTensor
        AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();

        // 打印调试信息（在NPU端）
        AscendC::printf("KernelAdd: 正在执行Add运算，数据长度=%u\n", TOTAL_LENGTH);

        // 执行Add计算：zLocal = xLocal + yLocal
        AscendC::Add(zLocal, xLocal, yLocal, TOTAL_LENGTH);

        // 打印完成信息
        AscendC::printf("KernelAdd: Add运算完成\n");

        // 将结果放入输出队列
        outQueueZ.EnQue<half>(zLocal);
        
        // 释放输入LocalTensor（归还给队列管理）
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    
    /**
     * CopyOut阶段：从本地内存拷贝结果回全局内存
     */
    __aicore__ inline void CopyOut()
    {
        // 从输出队列中取出结果LocalTensor
        AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
        
        // 从LocalTensor拷贝回GlobalTensor
        AscendC::DataCopy(zGm, zLocal, TOTAL_LENGTH);
        
        // 释放LocalTensor（归还给队列管理）
        outQueueZ.FreeTensor(zLocal);
    }

private:
    // TPipe用于管理内存和流水线
    AscendC::TPipe pipe;
    
    // TQue用于管理LocalTensor的分配和释放
    // TPosition::VECIN表示输入队列，VECOUT表示输出队列
    // 1表示队列的缓冲区数量
    AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueX;
    AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueY;
    AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueZ;
    
    // GlobalTensor用于访问全局内存
    AscendC::GlobalTensor<half> xGm;
    AscendC::GlobalTensor<half> yGm;
    AscendC::GlobalTensor<half> zGm;
};

/**
 * Add算子核函数
 * 
 * @param x 第一个输入张量的全局内存地址
 * @param y 第二个输入张量的全局内存地址
 * @param z 输出张量的全局内存地址
 */
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    AscendC::printf("add_custom核函数开始执行\n");
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
    AscendC::printf("add_custom核函数执行完成\n");
}

代码详解

1. GlobalTensor的创建和绑定

GlobalTensor<half> xGm, yGm, zGm;
xGm.SetGlobalBuffer((__gm__ half *)x, TOTAL_LENGTH);

• GlobalTensor<half>：创建half类型的全局张量对象
• SetGlobalBuffer：将全局内存地址绑定到GlobalTensor对象
• __gm__：全局内存修饰符，表示数据在全局内存中
• GM_ADDR：全局内存地址类型，由CPU端传入

2. LocalTensor的分配（重要：必须使用TPipe和TQue）

在Ascend C中，LocalTensor不能直接通过Alloc()方法分配。必须使用TPipe和TQue来管理：

// 1. 初始化TPipe和TQue
pipe.InitBuffer(inQueueX, 1, TOTAL_LENGTH * sizeof(half));

// 2. 从队列中分配LocalTensor
AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();

// 3. 使用完毕后，释放LocalTensor（归还给队列）
inQueueX.FreeTensor(xLocal);

• TPipe：用于管理内存和流水线
• TQue：队列，用于管理LocalTensor的分配和释放
• AllocTensor<half>()：从队列中分配LocalTensor（内存由TQue管理）
• FreeTensor()：释放LocalTensor，归还给队列
• 本地内存访问速度快，但容量有限（通常几百KB到几MB）

注意：这是Ascend C的标准做法，不能直接使用LocalTensor.Alloc()和LocalTensor.Free()。

3. 数据拷贝（CopyIn）

DataCopy(xLocal, xGm, TOTAL_LENGTH);

• DataCopy：Ascend C提供的数据拷贝API
• 从全局内存（xGm）拷贝到本地内存（xLocal）
• TOTAL_LENGTH：拷贝的元素个数
• 数据拷贝是异步的，可以使用流水线优化

4. Add计算

Add(zLocal, xLocal, yLocal, TOTAL_LENGTH);

• Add：Ascend C提供的向量加法API
• 函数签名：Add(dst, src1, src2, count)
• 执行：dst[i] = src1[i] + src2[i]，其中i = 0, 1, ..., count-1
• 内部使用向量指令，可以同时处理多个元素

5. 结果写回（CopyOut）

DataCopy(zGm, zLocal, TOTAL_LENGTH);

• 将计算结果从本地内存拷贝回全局内存
• CPU端可以从全局内存读取结果

6. 内存释放（使用TQue管理）

inQueueX.FreeTensor(xLocal);

• FreeTensor()：释放LocalTensor，归还给队列管理
• 必须显式释放，否则会导致内存泄漏
• 释放的顺序应该与分配的顺序相反

第二步：主程序实现（main.cpp）

修改main.cpp，添加数据准备、核函数调用和结果验证。注意：在CPU端代码中，不能使用__global__和__aicore__修饰符，这些只能在NPU端的核函数文件中使用。

/**
 * @file main.cpp
 * 
 * Add算子主程序 - 基于0_helloworld项目修改
 */
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>
#include <cstdint>

// 使用编译系统生成的头文件来调用核函数
// 这个头文件会在编译kernels库时自动生成
// 注意：需要先编译kernels库，然后才能编译main
#include "aclrtlaunch_add_custom.h"

// half类型在CPU端使用uint16_t表示（16位浮点数）
using half_t = uint16_t;

int32_t main(int argc, char const *argv[])
{
    printf("========================================\n");
    printf("Add算子测试 - 开始运行...\n");
    printf("========================================\n");
    
    // 1. 初始化ACL环境
    printf("步骤1: 初始化ACL环境...\n");
    aclInit(nullptr);
    int32_t deviceId = 0;
    aclrtSetDevice(deviceId);
    aclrtStream stream = nullptr;
    aclrtCreateStream(&stream);
    printf("  ACL环境初始化成功。\n");
    
    // 2. 数据长度
    constexpr uint32_t TOTAL_LENGTH = 2048;
    constexpr size_t dataSize = TOTAL_LENGTH * sizeof(half_t);
    printf("步骤2: 数据长度 = %u, 数据大小 = %zu 字节\n", TOTAL_LENGTH, dataSize);
    
    // 3. 准备Host端数据（简化：直接使用简单的值）
    printf("步骤3: 准备Host端数据...\n");
    half_t *host_x = (half_t *)malloc(dataSize);
    half_t *host_y = (half_t *)malloc(dataSize);
    half_t *host_z = (half_t *)malloc(dataSize);
    
    // 初始化简单的测试数据（使用简单的整数值，便于验证）
    for (uint32_t i = 0; i < TOTAL_LENGTH; i++) {
        // 使用简单的值：x[i] = i, y[i] = i*2, 期望结果 z[i] = i*3
        host_x[i] = (half_t)(i & 0xFFFF);
        host_y[i] = (half_t)((i * 2) & 0xFFFF);
    }
    printf("  Host端数据准备完成。前5个值：\n");
    for (uint32_t i = 0; i < 5; i++) {
        printf("    x[%u] = %u, y[%u] = %u\n", i, host_x[i], i, host_y[i]);
    }
    
    // 4. 在Device端分配全局内存
    printf("步骤4: 在Device端分配全局内存...\n");
    void *device_x = nullptr;
    void *device_y = nullptr;
    void *device_z = nullptr;
    
    aclrtMalloc(&device_x, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&device_y, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&device_z, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
    printf("  Device端内存分配完成。\n");
    
    // 5. 将数据从Host拷贝到Device
    printf("步骤5: 将数据从Host拷贝到Device...\n");
    aclrtMemcpy(device_x, dataSize, host_x, dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(device_y, dataSize, host_y, dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
    printf("  数据拷贝到Device完成。\n");
    
    // 6. 调用核函数
    printf("步骤6: 启动核函数...\n");
    constexpr uint32_t blockDim = 8;
    // 使用编译系统生成的宏来调用核函数
    ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, device_x, device_y, device_z);
    printf("  核函数已在 %u 个AI Core上启动。\n", blockDim);
    
    // 7. 同步等待核函数执行完成
    printf("步骤7: 同步等待核函数执行完成...\n");
    aclrtSynchronizeStream(stream);
    printf("  流同步完成，核函数执行完成。\n");
    
    // 8. 将结果从Device拷贝回Host
    printf("步骤8: 将结果从Device拷贝回Host...\n");
    aclrtMemcpy(host_z, dataSize, device_z, dataSize, ACL_MEMCPY_DEVICE_TO_HOST);
    printf("  结果拷贝到Host完成。\n");
    
    // 9. 打印结果
    printf("\n========================================\n");
    printf("计算结果：\n");
    printf("========================================\n");
    printf("前20个结果 (x + y = z)：\n");
    for (uint32_t i = 0; i < 20 && i < TOTAL_LENGTH; i++) {
        printf("  [%4u] %6u + %6u = %6u\n", 
               i, host_x[i], host_y[i], host_z[i]);
    }
    
    // 简单验证：检查前几个结果
    printf("\n验证结果（前10个元素）：\n");
    bool all_ok = true;
    for (uint32_t i = 0; i < 10 && i < TOTAL_LENGTH; i++) {
        uint32_t expected = host_x[i] + host_y[i];
        uint32_t got = host_z[i];
        if (expected != got) {
            printf("  [%u] 错误：期望值 %u，实际值 %u\n", i, expected, got);
            all_ok = false;
        } else {
            printf("  [%u] 正确：%u + %u = %u\n", i, host_x[i], host_y[i], got);
        }
    }
    
    printf("\n========================================\n");
    if (all_ok) {
        printf("测试通过！\n");
    } else {
        printf("测试失败！\n");
    }
    printf("========================================\n");
    
    // 10. 清理资源
    printf("\n步骤9: 清理资源...\n");
    free(host_x);
    free(host_y);
    free(host_z);
    aclrtFree(device_x);
    aclrtFree(device_y);
    aclrtFree(device_z);
    aclrtDestroyStream(stream);
    aclrtResetDevice(deviceId);
    aclFinalize();
    printf("  资源清理完成。\n");
    printf("========================================\n");
    
    return all_ok ? 0 : 1;
}

主程序详解

1. ACL环境初始化

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

• aclInit：初始化ACL（Ascend Computing Language）运行时环境
• aclrtSetDevice：设置使用的NPU设备
• aclrtCreateStream：创建异步执行流，用于异步执行核函数

2. 数据准备

half_t *host_x = (half_t *)malloc(dataSize);

• 在CPU端（Host）准备输入数据
• 使用half_t（即uint16_t）表示half类型，因为CPU端不能直接使用NPU端的half类型
• 使用malloc分配内存，使用完毕后用free释放
• 初始化简单的测试数据，便于验证结果

3. 设备内存分配

aclrtMalloc((void **)&device_x, TOTAL_LENGTH * sizeof(half), ACL_MEM_MALLOC_HUGE_FIRST);

• aclrtMalloc：在NPU设备端分配全局内存
• ACL_MEM_MALLOC_HUGE_FIRST：优先使用大页内存，提高性能

4. 数据拷贝

aclrtMemcpy(device_x, ..., host_x.data(), ..., ACL_MEMCPY_HOST_TO_DEVICE);

• aclrtMemcpy：在Host和Device之间拷贝数据
• ACL_MEMCPY_HOST_TO_DEVICE：从Host拷贝到Device
• ACL_MEMCPY_DEVICE_TO_HOST：从Device拷贝到Host

5. 核函数调用（使用ACLRT_LAUNCH_KERNEL宏）

ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, device_x, device_y, device_z);

• ACLRT_LAUNCH_KERNEL：编译系统自动生成的宏，用于调用核函数
• 这个宏定义在aclrtlaunch_add_custom.h头文件中
• blockDim = 8：在8个AI Core上并行执行
• 每个AI Core处理一部分数据
• 执行是异步的，不会阻塞CPU
• 注意：在CPU端代码中，不能直接声明核函数（不能使用__global__和__aicore__修饰符）

6. 结果验证

float error = std::abs(static_cast<float>(host_z[i] - host_z_golden[i]));

• 比较NPU计算结果和CPU计算结果
• 由于浮点数精度问题，允许一定的误差
• half精度下，误差通常在0.01以内

第三步：修改CMakeLists.txt

为了确保main.cpp能够找到编译系统生成的头文件aclrtlaunch_add_custom.h，需要修改CMakeLists.txt，添加包含路径：

add_executable(main main.cpp)

# 添加kernels库的头文件路径，以便main.cpp可以找到aclrtlaunch_add_custom.h
target_include_directories(main PRIVATE
    ${CMAKE_CURRENT_BINARY_DIR}/kernels_preprocess-prefix/src/kernels_preprocess-build/include
    ${CMAKE_INSTALL_PREFIX}/include
)

target_link_libraries(main PRIVATE
    kernels
)

说明：

• kernels_preprocess-build/include：编译时生成的头文件路径
• ${CMAKE_INSTALL_PREFIX}/include：安装后的头文件路径
• 这样main.cpp就可以找到aclrtlaunch_add_custom.h头文件了

Add算子的关键API详解

1. Add API

void Add(LocalTensor<DTYPE> &dst, 
         const LocalTensor<DTYPE> &src1, 
         const LocalTensor<DTYPE> &src2, 
         uint32_t count);

功能：执行向量加法运算

参数：

• dst：输出张量，存储计算结果
• src1：第一个输入张量
• src2：第二个输入张量
• count：参与计算的元素个数

执行：dst[i] = src1[i] + src2[i]，i = 0, 1, ..., count-1

支持的数据类型：half, float, int8_t, int16_t, int32_t等

性能特点：

• 使用向量指令，可以同时处理多个元素
• 对于half类型，通常可以同时处理256个元素
• 计算和内存访问可以流水线化

2. DataCopy API

void DataCopy(LocalTensor<DTYPE> &dst,
              const GlobalTensor<DTYPE> &src,
              uint32_t count);

功能：在全局内存和本地内存之间拷贝数据

参数：

• dst：目标张量（通常是LocalTensor）
• src：源张量（可以是GlobalTensor或LocalTensor）
• count：拷贝的元素个数

使用场景：

• CopyIn：从Global Memory拷贝到Local Memory
• CopyOut：从Local Memory拷贝到Global Memory
• 内部拷贝：在Local Memory之间拷贝

性能特点：

• 异步执行，可以与其他操作重叠
• 支持DMA（Direct Memory Access）加速
• 可以流水线化以提高效率

3. LocalTensor和GlobalTensor

LocalTensor

LocalTensor<half> xLocal;
xLocal.Alloc(count);      // 分配内存
xLocal.Free();            // 释放内存

特点：

• 存储在Unified Buffer（本地内存）中
• 访问速度快，延迟低
• 容量有限（通常几百KB到几MB）
• 用于临时存储和计算

GlobalTensor

GlobalTensor<half> xGm;
xGm.SetGlobalBuffer((__gm__ half *)addr, count);

特点：

• 存储在Global Memory（全局内存）中
• 容量大（通常几GB到几十GB）
• 访问速度相对较慢
• 用于存储输入输出数据

性能优化考虑

1. 数据分块（Tiling）

当数据量很大时，不能一次性将所有数据加载到Local Memory，需要分块处理：

constexpr uint32_t TILE_LENGTH = 256;  // 每块的大小

for (uint32_t i = 0; i < totalLength; i += TILE_LENGTH) {
    uint32_t currentLength = (i + TILE_LENGTH <= totalLength) ? 
                             TILE_LENGTH : (totalLength - i);
    
    // 加载当前块
    DataCopy(xLocal, xGm[i], currentLength);
    DataCopy(yLocal, yGm[i], currentLength);
    
    // 计算当前块
    Add(zLocal, xLocal, yLocal, currentLength);
    
    // 写回当前块
    DataCopy(zGm[i], zLocal, currentLength);
}

2. 流水线优化

使用TPipe和TQue实现流水线，让CopyIn、Compute、CopyOut三个阶段重叠执行：

class KernelAdd {
private:
    TPipe pipe;
    TQue<TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
    TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ;
    
public:
    void Process() {
        // 流水线处理多个块
        for (uint32_t i = 0; i < tileNum; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }
};

3. 多核并行

通过blockDim参数控制使用的AI Core数量：

constexpr uint32_t blockDim = 8;  // 使用8个AI Core

每个AI Core处理一部分数据，实现数据并行。

完整示例代码

hello_world.cpp（完整版）

完整的核函数实现代码请参考上面的"第一步：核函数实现"部分，这里不再重复。关键点：

1. 使用KernelAdd类封装逻辑
2. 使用TPipe和TQue管理LocalTensor
3. 通过TQue.AllocTensor()分配LocalTensor
4. 通过TQue.FreeTensor()释放LocalTensor
5. 使用AscendC::printf()打印调试信息（中文）

main.cpp（完整版）

完整的main.cpp代码请参考上面的"第二步：主程序实现"部分。关键点：

1. 包含aclrtlaunch_add_custom.h头文件（编译系统自动生成）
2. 使用ACLRT_LAUNCH_KERNEL宏调用核函数
3. 使用half_t（即uint16_t）表示half类型
4. 使用printf打印中文提示信息
5. 使用malloc/free管理Host端内存

基于0_helloworld项目的修改步骤

步骤1：修改hello_world.cpp

直接修改hello_world.cpp文件，实现Add算子。完整的代码请参考上面的"第一步：核函数实现"部分。

关键点：

• 使用KernelAdd类封装逻辑
• 使用TPipe和TQue管理LocalTensor
• 通过TQue.AllocTensor()分配LocalTensor，不能直接使用Alloc()
• 通过TQue.FreeTensor()释放LocalTensor，不能直接使用Free()
• 使用AscendC::printf()打印中文调试信息

步骤2：修改main.cpp

修改main.cpp，添加数据准备和验证逻辑。完整的代码请参考上面的"第二步：主程序实现"部分。

关键点：

• 包含aclrtlaunch_add_custom.h头文件（编译系统自动生成）
• 使用ACLRT_LAUNCH_KERNEL宏调用核函数
• 使用half_t（即uint16_t）表示half类型
• 使用printf打印中文提示信息
• 注意：在CPU端代码中，不能使用__global__和__aicore__修饰符

步骤3：修改CMakeLists.txt

修改CMakeLists.txt，添加头文件包含路径：

add_executable(main main.cpp)

# 添加kernels库的头文件路径，以便main.cpp可以找到aclrtlaunch_add_custom.h
target_include_directories(main PRIVATE
    ${CMAKE_CURRENT_BINARY_DIR}/kernels_preprocess-prefix/src/kernels_preprocess-build/include
    ${CMAKE_INSTALL_PREFIX}/include
)

target_link_libraries(main PRIVATE
    kernels
)

说明：

• hello_world.cpp文件名保持不变（CMakeLists.txt中也是hello_world.cpp）
• 添加target_include_directories确保能找到编译生成的头文件

步骤4：编译和运行

cd samples/operator/ascendc/0_introduction/0_helloworld
bash run.sh -v Ascend910B4

注意：根据你的实际NPU型号修改-v参数。可以通过npu-smi info命令查看NPU型号。

预期输出

如果运行成功，应该看到类似以下的中文输出：

========================================
Add算子测试 - 开始运行...
========================================
步骤1: 初始化ACL环境...
  ACL环境初始化成功。
步骤2: 数据长度 = 2048, 数据大小 = 4096 字节
步骤3: 准备Host端数据...
  Host端数据准备完成。前5个值：
    x[0] = 0, y[0] = 0
    x[1] = 1, y[1] = 2
    ...
步骤4: 在Device端分配全局内存...
  Device端内存分配完成。
...
========================================
计算结果：
========================================
前20个结果 (x + y = z)：
  [   0]      0 +      0 =      0
  [   1]      1 +      2 =      3
  ...

验证结果（前10个元素）：
  [0] 正确：0 + 0 = 0
  [1] 正确：1 + 2 = 3
  ...

========================================
测试通过！
========================================

同时，在NPU端也会打印调试信息（通过AscendC::printf）：

add_custom核函数开始执行
KernelAdd: 正在执行Add运算，数据长度=2048
KernelAdd: Add运算完成
add_custom核函数执行完成

关键注意事项

1. LocalTensor的正确使用方式

重要：在Ascend C中，LocalTensor不能直接使用Alloc()和Free()方法。必须使用TPipe和TQue来管理：

// ❌ 错误的方式（会导致编译错误）
LocalTensor<half> xLocal;
xLocal.Alloc(TOTAL_LENGTH);  // 编译错误：没有Alloc方法
xLocal.Free();               // 编译错误：没有Free方法

// ✅ 正确的方式
TPipe pipe;
TQue<TPosition::VECIN, 1> inQueueX;
pipe.InitBuffer(inQueueX, 1, TOTAL_LENGTH * sizeof(half));
LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();  // 正确
inQueueX.FreeTensor(xLocal);  // 正确

2. CPU端和NPU端的代码分离

• NPU端代码（hello_world.cpp）：可以使用__global__、__aicore__、GM_ADDR、half等NPU端类型
• CPU端代码（main.cpp）：不能使用NPU端修饰符和类型，需要使用标准C++类型

3. 核函数调用方式

使用编译系统自动生成的ACLRT_LAUNCH_KERNEL宏，而不是直接调用：

// ✅ 正确的方式
#include "aclrtlaunch_add_custom.h"
ACLRT_LAUNCH_KERNEL(add_custom)(blockDim, stream, device_x, device_y, device_z);

// ❌ 错误的方式（在CPU端不能这样声明）
extern "C" __global__ __aicore__ void add_custom(...);  // 编译错误

4. 数据类型转换

在CPU端，half类型需要使用uint16_t表示：

// CPU端
using half_t = uint16_t;
half_t *host_x = (half_t *)malloc(dataSize);

// NPU端
half *xGm;  // 直接使用half类型

总结

本文详细介绍了Add算子的实现，包括：

1. 正确的实现方式：使用TPipe和TQue管理LocalTensor
2. API详解：Add、DataCopy等关键API的使用方法
3. 内存管理：Global Memory和Local Memory的区别和使用
4. 代码分离：CPU端和NPU端代码的正确编写方式
5. 编译配置：CMakeLists.txt的正确配置

Add算子虽然简单，但它包含了Ascend C算子开发的核心要素。掌握了Add算子的实现，就为学习更复杂的算子打下了坚实的基础。

在下一篇文章中，我们将介绍Sub、Mul、Div等其他元素级算子的实现，它们与Add算子的实现方式非常相似，主要区别在于使用的API不同。

---

2025年昇腾CANN训练营第二季，基于CANN开源开放全场景，推出0基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机，平板、开发板等大奖。

报名链接:https://www.hiascend.com/developer/activities/cann20252

社区地址：https://www.hiascend.com/developer