概述

Div算子（Division Operator）是元素级算子的一种，用于实现两个张量的逐元素相除。Div算子与Add、Sub、Mul算子非常相似，主要区别在于使用的API不同：Add使用Add API，Sub使用Sub API，Mul使用Mul API，而Div使用Div API。本文将在Add、Sub、Mul算子的基础上，展示如何实现Div算子，并同步更新0_helloworld项目的代码。

什么是Div算子

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

output[i] = input1[i] / input2[i]

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

Div算子的特点

1. 元素独立性：每个输出元素只依赖于对应位置的输入元素，元素之间没有依赖关系
2. 易于并行化：由于元素独立性，可以充分利用多核并行计算
3. 易于向量化：可以使用向量指令同时处理多个元素
4. 内存访问模式简单：顺序访问，缓存友好
5. 需要注意除零问题：当除数为0时，会产生无穷大或NaN，需要在应用层处理

Div算子的应用场景

• 归一化操作：将张量除以某个归一化因子
• 特征缩放：对特征进行缩放处理
• 注意力机制：在Transformer等模型中，用于计算注意力权重
• 梯度更新：在优化算法中，用于更新参数
• 广播除法：支持不同形状张量的广播相除

基于0_helloworld实现Div算子

我们将基于0_helloworld项目来实现Div算子。由于Div算子与Add、Sub、Mul算子非常相似，我们只需要将API替换为Div API即可。

项目结构

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

0_helloworld/
├── CMakeLists.txt          # 编译配置文件（基本不变）
├── hello_world.cpp         # 修改核函数实现（Mul改为Div）
├── main.cpp                # 修改主程序（Mul改为Div，更新验证逻辑）
└── run.sh                  # 运行脚本（基本不变）

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

Div算子的实现与Add、Sub、Mul算子几乎完全相同，只需要将Mul API替换为Div API：

/**
 * @file hello_world.cpp
 * 
 * Div算子实现 - 基于0_helloworld项目修改
 * 对应第二十二篇：Div算子实现详解
 * 
 * Copyright (c) 2025
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
 */
#include "kernel_operator.h"

constexpr int32_t TOTAL_LENGTH = 8 * 2048;                            // total length of data
constexpr int32_t USE_CORE_NUM = 8;                                   // num of core used
constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM;         // length computed of each core
constexpr int32_t TILE_NUM = 8;                                       // split data into 8 tiles for each core
constexpr int32_t BUFFER_NUM = 2;                                     // tensor num for each queue
constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // separate to 2 parts, due to double buffer

/**
 * Div算子Kernel类
 * 使用TPipe和TQue来管理LocalTensor的内存分配
 */
class KernelDiv {
public:
    __aicore__ inline KernelDiv() {}
    
    /**
     * 初始化函数
     * @param x 第一个输入张量的全局内存地址（被除数）
     * @param y 第二个输入张量的全局内存地址（除数）
     * @param z 输出张量的全局内存地址（商）
     */
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
    {
        // 1. 创建GlobalTensor对象，绑定全局内存
        // 每个Core只处理一部分数据，使用GetBlockIdx()获取当前Core的索引
        int32_t blockIdx = AscendC::GetBlockIdx();
        xGm.SetGlobalBuffer((__gm__ half *)x + BLOCK_LENGTH * blockIdx, BLOCK_LENGTH);
        yGm.SetGlobalBuffer((__gm__ half *)y + BLOCK_LENGTH * blockIdx, BLOCK_LENGTH);
        zGm.SetGlobalBuffer((__gm__ half *)z + BLOCK_LENGTH * blockIdx, BLOCK_LENGTH);

        // 添加调试信息（只在第一个Core打印，避免输出过多）
        if (blockIdx == 0) {
            AscendC::printf("KernelDiv::Init: blockIdx=%d, BLOCK_LENGTH=%d, TILE_LENGTH=%d\n", 
                            blockIdx, BLOCK_LENGTH, TILE_LENGTH);
        }

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

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

        // 从GlobalTensor拷贝到LocalTensor（使用偏移量）
        AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
        AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);

        // 将LocalTensor放入队列（用于后续的Compute阶段）
        inQueueX.EnQue(xLocal);
        inQueueY.EnQue(yLocal);
        
        // 添加调试信息（只在第一个Core的第一个tile打印）
        if (AscendC::GetBlockIdx() == 0 && progress == 0) {
            AscendC::printf("KernelDiv::CopyIn: progress=%d, TILE_LENGTH=%d完成\n", progress, TILE_LENGTH);
        }
    }
    
    /**
     * Compute阶段：执行Div计算
     */
    __aicore__ inline void Compute(int32_t progress)
    {
        // 从队列中取出LocalTensor
        AscendC::LocalTensor<half> xLocal = inQueueX.DeQue<half>();
        AscendC::LocalTensor<half> yLocal = inQueueY.DeQue<half>();
        
        // 为输出分配LocalTensor
        AscendC::LocalTensor<half> zLocal = outQueueZ.AllocTensor<half>();

        // 添加调试信息（只在第一个tile打印，避免输出过多）
        if (progress == 0) {
            AscendC::printf("KernelDiv: 开始执行Div计算, progress=%d, TILE_LENGTH=%d\n", progress, TILE_LENGTH);
        }

        // 执行Div计算：zLocal = xLocal / yLocal
        AscendC::Div(zLocal, xLocal, yLocal, TILE_LENGTH);

        if (progress == 0) {
            AscendC::printf("KernelDiv: Div计算完成, progress=%d\n", progress);
        }

        // 将结果放入输出队列
        outQueueZ.EnQue<half>(zLocal);
        
        // 释放输入LocalTensor（归还给队列管理）
        inQueueX.FreeTensor(xLocal);
        inQueueY.FreeTensor(yLocal);
    }
    
    /**
     * CopyOut阶段：从本地内存拷贝结果回全局内存
     */
    __aicore__ inline void CopyOut(int32_t progress)
    {
        // 从输出队列中取出结果LocalTensor
        AscendC::LocalTensor<half> zLocal = outQueueZ.DeQue<half>();
        
        // 从LocalTensor拷贝回GlobalTensor（使用偏移量）
        AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
        
        // 添加调试信息（只在第一个Core的第一个tile打印）
        if (AscendC::GetBlockIdx() == 0 && progress == 0) {
            AscendC::printf("KernelDiv::CopyOut: progress=%d, TILE_LENGTH=%d完成\n", progress, TILE_LENGTH);
        }
        
        // 释放LocalTensor（归还给队列管理）
        outQueueZ.FreeTensor(zLocal);
    }

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

/**
 * Div算子核函数
 * 
 * @param x 第一个输入张量的全局内存地址（被除数）
 * @param y 第二个输入张量的全局内存地址（除数）
 * @param z 输出张量的全局内存地址（商）
 */
extern "C" __global__ __aicore__ void div_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    // 添加调试信息，确认核函数被调用
    AscendC::printf("div_custom: 核函数开始执行, blockIdx=%d\n", AscendC::GetBlockIdx());
    KernelDiv op;
    op.Init(x, y, z);
    AscendC::printf("div_custom: Init完成, blockIdx=%d\n", AscendC::GetBlockIdx());
    op.Process();
    AscendC::printf("div_custom: Process完成, blockIdx=%d\n", AscendC::GetBlockIdx());
}

#ifndef ASCENDC_CPU_DEBUG
/**
 * 核函数调用包装函数
 * 用于从CPU端调用NPU核函数
 * 注意：参数类型使用uint8_t*，与Add算子保持一致
 */
void div_custom_do(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, uint8_t *z)
{
    div_custom<<<blockDim, nullptr, stream>>>(x, y, z);
}
#endif

代码详解

关键修改点：Mul API → Div API

与Mul算子相比，Div算子的唯一区别在于Compute阶段使用的API：

// Mul算子
AscendC::Mul(zLocal, xLocal, yLocal, TILE_LENGTH);  // zLocal = xLocal * yLocal

// Div算子
AscendC::Div(zLocal, xLocal, yLocal, TILE_LENGTH);  // zLocal = xLocal / yLocal

Div API详解

void Div(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等浮点类型

性能特点：

• 使用向量指令，可以同时处理多个元素
• 对于half类型，通常可以同时处理256个元素
• 计算和内存访问可以流水线化
• 与Add、Sub、Mul API性能相当

注意事项：

• 当除数为0时，会产生无穷大（Inf）或NaN（Not a Number）
• 在实际应用中，需要检查除数是否为0，避免除零错误
• 对于整数类型，除法结果会向下取整

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

修改main.cpp，将Mul改为Div，并更新数据验证逻辑：

/**
 * @file main.cpp
 * 
 * Div算子主程序 - 基于0_helloworld项目修改
 * 对应第二十二篇：Div算子实现详解
 * 
 * Copyright (c) 2025
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
 */
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>
#include <cstdint>

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

// 声明包装函数
#ifndef ASCENDC_CPU_DEBUG
extern void div_custom_do(uint32_t blockDim, void *stream, uint8_t *x, uint8_t *y, uint8_t *z);
#endif

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

int32_t main(int argc, char const *argv[])
{
    printf("========================================\n");
    printf("Div算子测试 - 开始运行...\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. 数据长度（与核函数中的TOTAL_LENGTH保持一致）
    constexpr uint32_t TOTAL_LENGTH = 8 * 2048;  // 8 * 2048 = 16384
    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);
    
    // 初始化简单的测试数据
    // 注意：half是16位浮点数，不能直接使用整数赋值
    // 这里我们使用简单的位模式，但实际half类型需要正确的浮点编码
    // Div: x[i] / y[i] = z[i]
    // 使用简单的值：x[i] = i*2, y[i] = 2, 期望结果 z[i] = i
    // 注意：确保除数y[i]不为0，避免除零错误
    // 为了简化，我们直接使用整数位模式（虽然不正确，但可以验证计算流程）
    for (uint32_t i = 0; i < TOTAL_LENGTH; i++) {
        // 直接使用整数位模式（注意：这不是正确的half浮点数编码）
        host_x[i] = (half_t)(i * 2);  // 被除数（确保不为0）
        host_y[i] = (half_t)(2);      // 除数（固定为2，确保不为0）
    }
    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);
    
    // 初始化device_z为0（确保输出内存是干净的）
    half_t *zero_data = (half_t *)calloc(TOTAL_LENGTH, sizeof(half_t));
    aclrtMemcpy(device_z, dataSize, zero_data, dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
    free(zero_data);
    
    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;
    // 使用编译系统生成的宏来调用核函数
    printf("  准备调用核函数 div_custom，参数：blockDim=%u, stream=%p\n", blockDim, stream);
    printf("  device_x=%p, device_y=%p, device_z=%p\n", device_x, device_y, device_z);
    
    // 直接使用包装函数调用核函数
    // 这样可以避免ACLRT_LAUNCH_KERNEL宏的问题
    #ifndef ASCENDC_CPU_DEBUG
        printf("  调用div_custom_do函数...\n");
        fflush(stdout);  // 确保输出立即刷新
        div_custom_do(blockDim, stream, (uint8_t *)device_x, (uint8_t *)device_y, (uint8_t *)device_z);
        printf("  div_custom_do函数调用完成。\n");
        fflush(stdout);
    #else
        // CPU调试模式使用不同的调用方式
        printf("  错误：不支持CPU调试模式\n");
        return 1;
    #endif
    
    printf("  核函数已在 %u 个AI Core上启动。\n", blockDim);
    
    // 7. 同步等待核函数执行完成
    printf("步骤7: 同步等待核函数执行完成...\n");
    aclError ret = aclrtSynchronizeStream(stream);
    if (ret != ACL_SUCCESS) {
        printf("  错误：流同步失败，错误码=%d\n", ret);
    } else {
        printf("  流同步完成，核函数执行完成。\n");
    }
    
    // 8. 将结果从Device拷贝回Host
    printf("步骤8: 将结果从Device拷贝回Host...\n");
    aclError memcpy_ret = aclrtMemcpy(host_z, dataSize, device_z, dataSize, ACL_MEMCPY_DEVICE_TO_HOST);
    if (memcpy_ret != ACL_SUCCESS) {
        printf("  错误：结果拷贝失败，错误码=%d\n", memcpy_ret);
    } else {
        printf("  结果拷贝到Host完成。\n");
    }
    
    // 打印一些调试信息
    printf("\n========================================\n");
    printf("详细调试信息：\n");
    printf("========================================\n");
    printf("数据长度: TOTAL_LENGTH=%u, dataSize=%zu字节\n", TOTAL_LENGTH, dataSize);
    printf("核函数参数: blockDim=8, 每个Core处理BLOCK_LENGTH=2048个元素\n");
    printf("前10个输入值（Host端）：\n");
    for (uint32_t i = 0; i < 10 && i < TOTAL_LENGTH; i++) {
        printf("    x[%u] = 0x%04X (%u), y[%u] = 0x%04X (%u)\n", 
               i, host_x[i], host_x[i], i, host_y[i], host_y[i]);
    }
    printf("前10个输出值（Host端，从Device拷贝后）：\n");
    for (uint32_t i = 0; i < 10 && i < TOTAL_LENGTH; i++) {
        printf("    z[%u] = 0x%04X (%u)\n", i, host_z[i], host_z[i]);
    }
    
    // 检查是否有非零值
    uint32_t non_zero_count = 0;
    for (uint32_t i = 0; i < TOTAL_LENGTH; i++) {
        if (host_z[i] != 0) {
            non_zero_count++;
            if (non_zero_count <= 5) {
                printf("  发现非零值: z[%u] = 0x%04X (%u)\n", i, host_z[i], host_z[i]);
            }
        }
    }
    printf("非零值统计: 共 %u/%u 个非零值\n", non_zero_count, TOTAL_LENGTH);
    printf("========================================\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]);
    }
    
    // 简单验证：检查前几个结果
    // 注意：由于我们使用的是整数位模式而不是真正的half浮点数，
    // 所以验证时也使用整数位模式比较
    printf("\n验证结果（前10个元素）：\n");
    bool all_ok = true;
    for (uint32_t i = 0; i < 10 && i < TOTAL_LENGTH; i++) {
        // 使用整数位模式进行比较（因为输入也是整数位模式）
        // Div: x / y
        // 注意：整数除法会向下取整
        uint32_t expected = (uint32_t)host_x[i] / (uint32_t)host_y[i];  // Div: x / y
        uint32_t got = (uint32_t)host_z[i];
        if (expected != got) {
            printf("  [%u] 错误：期望值 %u，实际值 %u (x=0x%04X, y=0x%04X, z=0x%04X)\n", 
                   i, expected, got, host_x[i], host_y[i], host_z[i]);
            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;
}

代码详解

关键修改点：Mul → Div

与Mul算子相比，Div算子的主要修改点：

1. 核函数名称：mul_custom → div_custom
2. 类名：KernelMul → KernelDiv
3. API调用：AscendC::Mul → AscendC::Div
4. 数据初始化：x[i] = i, y[i] = 2 → x[i] = i*2, y[i] = 2（确保被除数不为0）
5. 验证逻辑：x * y → x / y

除零处理

在实际应用中，需要特别注意除零问题：

// 在实际应用中，可以添加除零检查
if (host_y[i] == 0) {
    // 处理除零情况：可以设置为0、NaN或跳过
    host_z[i] = 0;  // 或其他处理方式
} else {
    host_z[i] = host_x[i] / host_y[i];
}

但在Ascend C的Div API中，当除数为0时，会自动产生无穷大（Inf）或NaN（Not a Number），这符合IEEE 754浮点数标准。

总结

Div算子的实现与Add、Sub、Mul算子非常相似，主要区别在于：

1. API不同：使用Div API而不是Add、Sub或Mul API
2. 数学运算不同：执行除法运算而不是加法、减法或乘法运算
3. 需要注意除零问题：在实际应用中，需要检查除数是否为0

Div算子的实现同样遵循了Ascend C的标准模式：

• 使用TPipe和TQue管理LocalTensor的内存
• 使用分块（Tiling）和多核并行处理
• 使用双缓冲（Double Buffering）优化性能
• 使用流水线（Pipeline）重叠数据传输和计算

通过本文的学习，读者应该能够：

1. 理解Div算子的基本概念和应用场景
2. 掌握如何使用Div API实现逐元素除法
3. 了解Div算子与Add、Sub、Mul算子的异同
4. 理解除零问题的处理方式

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

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

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