我第一次没测试成功：

二次测试，是当前的下面的代码哦，这回肯定没问题：

概述

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

实际演示效果：

什么是Mul算子

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

output[i] = input1[i] * input2[i]

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

Mul算子的特点

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

Mul算子的应用场景

• 特征融合：将两个特征图逐元素相乘，实现特征融合
• 注意力机制：在Transformer等模型中，用于计算注意力权重
• 门控机制：在LSTM、GRU等模型中，用于实现门控功能
• 广播乘法：支持不同形状张量的广播相乘
• 缩放操作：对张量进行逐元素缩放

实现Mul算子

由于Mul算子与Add、Sub算子非常相似，我们只需要将API替换为Mul API即可。

项目结构

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

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

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

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

/**
 * @file hello_world.cpp
 * 
 * Mul算子实现 - 基于0_helloworld项目修改
 * 对应第二十一篇：Mul算子实现详解
 * 
 * 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 uint32_t TOTAL_LENGTH = 2048;

/**
 * Mul算子Kernel类
 * 使用TPipe和TQue来管理LocalTensor的内存分配
 */
class KernelMul {
public:
    __aicore__ inline KernelMul() {}
    
    /**
     * 初始化函数
     * @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));
    }
    
    /**
     * 处理函数，执行完整的Mul算子流程
     */
    __aicore__ inline void Process()
    {
        CopyIn();   // 从全局内存拷贝到本地内存
        Compute();  // 执行Mul计算
        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阶段：执行Mul计算
     */
    __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("KernelMul: 正在执行Mul运算，数据长度=%u\n", TOTAL_LENGTH);

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

        // 打印完成信息
        AscendC::printf("KernelMul: Mul运算完成\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;
};

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

代码详解

关键修改点：Sub API → Mul API

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

// Sub算子
AscendC::Sub(zLocal, xLocal, yLocal, TOTAL_LENGTH);  // zLocal = xLocal - yLocal

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

Mul API详解

void Mul(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个元素
• 计算和内存访问可以流水线化
• 与Add、Sub API性能相当

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

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

/**
 * @file main.cpp
 * 
 * Mul算子主程序 - 基于0_helloworld项目修改
 * 对应第二十一篇：Mul算子实现详解
 * 
 * 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_mul_custom.h"

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

int32_t main(int argc, char const *argv[])
{
    printf("========================================\n");
    printf("Mul算子测试 - 开始运行...\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);
    
    // 初始化简单的测试数据（使用简单的整数值，便于验证）
    // Mul: x[i] * y[i] = z[i]
    // 使用简单的值：x[i] = i, y[i] = 2, 期望结果 z[i] = i*2
    for (uint32_t i = 0; i < TOTAL_LENGTH; i++) {
        host_x[i] = (half_t)(i & 0xFFFF);  // 被乘数
        host_y[i] = (half_t)(2 & 0xFFFF);  // 乘数（固定为2）
    }
    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(mul_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];  // Mul: x * y
        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. 头文件：从aclrtlaunch_sub_custom.h改为aclrtlaunch_mul_custom.h
2. 核函数调用：从ACLRT_LAUNCH_KERNEL(sub_custom)改为ACLRT_LAUNCH_KERNEL(mul_custom)
3. 数据初始化：更新为Mul算子的测试数据模式（x[i] = i, y[i] = 2）
4. 验证逻辑：从x - y改为x * y
5. 时间备注：添加了Copyright (c) 2025时间备注

Mul算子与Add、Sub算子的对比

相同点

1. 代码结构完全相同：都使用KernelAdd/KernelSub/KernelMul类，都使用TPipe和TQue管理内存
2. 内存管理方式相同：都使用相同的LocalTensor分配和释放方式
3. 数据流相同：CopyIn → Compute → CopyOut三个阶段
4. 调用方式相同：都使用ACLRT_LAUNCH_KERNEL宏

不同点

1. API不同：

   • Add使用AscendC::Add()
   • Sub使用AscendC::Sub()
   • Mul使用AscendC::Mul()

2. 数学运算不同：

   • Add: z[i] = x[i] + y[i]
   • Sub: z[i] = x[i] - y[i]
   • Mul: z[i] = x[i] * y[i]

3. 测试数据不同：

   • Add: x[i] = i, y[i] = i*2, z[i] = i*3
   • Sub: x[i] = i*3, y[i] = i, z[i] = i*2
   • Mul: x[i] = i, y[i] = 2, z[i] = i*2

Mul API详解

Mul API函数签名

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

参数说明

• dst：输出张量，存储计算结果 dst[i] = src1[i] * src2[i]
• src1：第一个输入张量（被乘数）
• src2：第二个输入张量（乘数）
• count：参与计算的元素个数

支持的数据类型

• 浮点类型：half, float
• 整数类型：int8_t, int16_t, int32_t

性能特点

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

相关API：Muls（标量乘法）

除了Mul API，Ascend C还提供了Muls API，用于向量与标量的乘法：

void Muls(LocalTensor<DTYPE> &dst, 
          const LocalTensor<DTYPE> &src, 
          const DTYPE scalar, 
          uint32_t count);

功能：dst[i] = src[i] * scalar，即向量每个元素乘以标量

使用场景：当需要将整个向量乘以一个常数时，使用Muls比Mul更高效

基于0_helloworld项目的修改步骤

步骤1：修改hello_world.cpp

将Sub算子改为Mul算子：

1. 将KernelSub类名改为KernelMul
2. 将Sub API改为Mul API
3. 将sub_custom函数名改为mul_custom
4. 更新printf中的提示信息（Sub改为Mul）
5. 添加Copyright (c) 2025时间备注

步骤2：修改main.cpp

1. 将头文件从aclrtlaunch_sub_custom.h改为aclrtlaunch_mul_custom.h
2. 将ACLRT_LAUNCH_KERNEL(sub_custom)改为ACLRT_LAUNCH_KERNEL(mul_custom)
3. 更新数据初始化逻辑（改为Mul的测试数据：x[i] = i, y[i] = 2）
4. 更新验证逻辑（从x - y改为x * y）
5. 更新所有printf中的提示信息（Sub改为Mul）
6. 添加Copyright (c) 2025时间备注

步骤3：编译和运行

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

预期输出

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

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

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

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

同时，在NPU端也会打印调试信息：

mul_custom核函数开始执行
KernelMul: 正在执行Mul运算，数据长度=2048
KernelMul: Mul运算完成
mul_custom核函数执行完成

关键注意事项

1. API的选择

Mul算子使用Mul API，函数签名与Add、Sub API完全相同，只是运算不同：

// Add: z = x + y
AscendC::Add(zLocal, xLocal, yLocal, TOTAL_LENGTH);

// Sub: z = x - y
AscendC::Sub(zLocal, xLocal, yLocal, TOTAL_LENGTH);

// Mul: z = x * y
AscendC::Mul(zLocal, xLocal, yLocal, TOTAL_LENGTH);

2. 乘法运算的特点

• 乘法运算比加法和减法运算稍慢，但Ascend C的向量化实现仍然非常高效
• 注意数值溢出问题，特别是整数类型
• 对于浮点类型，注意精度问题

3. 与标量乘法的区别

• Mul：向量与向量逐元素相乘
• Muls：向量与标量相乘（更高效）

4. 时间备注

代码中添加了Copyright (c) 2025时间备注，表明这是2025年的项目代码。

总结

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

1. Mul算子的定义：逐元素相乘运算
2. 实现方式：使用TPipe和TQue管理LocalTensor
3. API详解：Mul API的使用方法和参数说明
4. 与Add、Sub算子的对比：相同点和不同点
5. 代码修改步骤：如何从Sub算子改为Mul算子
6. 时间备注：添加了2025年的版权时间备注

Mul算子与Add、Sub算子非常相似，掌握了Add和Sub算子的实现后，实现Mul算子就非常简单了，只需要替换API即可。这体现了Ascend C算子开发的模块化和可复用性。

在下一篇文章中，我们将介绍Div（除法）算子的实现，完成四则运算算子的完整系列。

---

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

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

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