华为ascend c 示例代码和详细讲解，昇腾 CANN 算子双版本模板代码

华为ascend c 示例代码和详细讲解，昇腾 CANN 算子双版本模板代码

下面给逐行详细注释 + 整体逻辑讲解，基于昇腾CANN 双分支（CCE仿真测试 / ACL 设备运行）场景，同时补充昇腾相关概念。

ai生成

// ==============================================
// 宏分支区分编译环境：CCE内核仿真测试 / 标准ACL应用
// ==============================================
// 如果定义了 CCE_KT_TEST：走CCE内核态本地仿真（CPU模拟aicore运行，单元测试用）
#ifdef CCE_KT_TEST
// 引入TIK/CCE仿真库头文件，提供内存分配、内核调用等仿真接口
#include "tikicpulib.h"
// 定义aicore为空宏，适配aicore内核函数修饰符语法
#define aicore
// 未定义 CCE_KT_TEST：走正式ACL应用层，运行在真实昇腾AI芯片上
#else
// 引入昇腾ACL基础接口头文件（设备、流、内存、上下文管理）
#include "acl/acl.h"
// 同样定义aicore空宏，统一两端代码语法
#define aicore
#endif

// ==============================================
// AICORE 内核函数声明（算子/核函数入口）
// ==============================================
// extern "C"：禁用C++名字修饰，保证C/CCE框架能正常调用该函数
// aicore_void：aicore内核侧空返回值类型（由上面空宏兼容）
// HelloWorld：自定义aicore核函数名
// uint8_t* foo：核函数入参，指向设备全局内存(Gm)数据
extern "C" aicore_void HelloWorld(uint8_t* foo);

// 主函数入口，标准程序入口
int32_t main(int32_t argc, char* argv[])
{
// 定义内存大小：申请 256 字节的全局内存空间
size_t fooSize = 256;
// 定义执行块维度：一次启动 8 个aicore执行单元并行计算
uint32_t blockDim = 8;

// ==============================================
// 分支1：CCE_KT_TEST 内核仿真模式（本地CPU调试）
// ==============================================
#ifdef CCE_KT_TEST
// 调用CCE仿真库接口，分配芯片全局内存(Gm)，返回内存指针
// Ascendl::GmAlloc：仿真环境下的全局内存分配接口
uint8_t* foo = (uint8_t*)Ascendl::GmAlloc(fooSize);

// ICPU_RUN_KF：CCE仿真框架调用aicore内核函数的宏
// 参数1：核函数名 HelloWorld
// 参数2：并行执行单元数 blockDim
// 参数3：传给核函数的内存指针 foo
ICPU_RUN_KF(HelloWorld, blockDim, foo);

// 释放之前分配的全局内存，避免内存泄漏
Ascendl::GmFree((void *)foo);

// ==============================================
// 分支2：标准ACL模式（真实昇腾硬件运行，业务正式环境）
// ==============================================
#else
// 1. 初始化ACL框架，整个程序生命周期只调用一次
// nullptr：使用默认初始化参数
aclInit(nullptr);

// 定义ACL上下文句柄：上下文用于管理设备、流、任务等资源
aclrtContext context;
// 指定使用的昇腾设备ID，此处使用第0号设备
int32_t deviceId = 0;
// 绑定当前线程到指定昇腾设备
aclrtSetDevice(deviceId);
// 基于设备ID创建运行上下文
aclrtCreateContext(&context, deviceId);

// 定义流句柄：Stream 用于异步调度AI任务、控制任务执行顺序
aclrtStream stream = nullptr;
// 创建一个空的计算流
aclrtCreateStream(&stream);

// 定义设备侧内存指针，初始置空
uint8_t* fooDevice = nullptr;

// 在昇腾设备上分配显存/全局内存
// &fooDevice：接收内存指针
// fooSize：内存大小
// ACL_MEM_MALLOC_HUGE_FIRST：优先使用大页内存，提升访存性能
aclrtMalloc((void**)&fooDevice, fooSize, ACL_MEM_MALLOC_HUGE_FIRST);

// 昇腾核函数调用语法 <<<块数, 共享内存大小, 流>>>
// blockDim：并行执行单元数
// nullptr：不使用共享内存
// stream：任务挂载到指定流上异步执行
// 后面括号：核函数入参（设备内存指针）
HelloWorld<<<blockDim, nullptr, stream>>>(fooDevice);

// 流同步：阻塞当前线程，等待流上所有任务执行完毕
// 保证核函数跑完再执行后续逻辑
aclrtSynchronizeStream(stream);

// 逆序释放资源（创建顺序：内存→流→上下文→设备→ACL，释放反过来）
// 释放设备侧内存
aclrtFree(fooDevice);
// 销毁计算流
aclrtDestroyStream(stream);
// 销毁运行上下文
aclrtDestroyContext(context);
// 解除当前线程与设备的绑定
aclrtResetDevice(deviceId);
// 最终销毁ACL框架，程序退出前最后一步
aclFinalize();
#endif

// 程序正常退出，返回0
return 0;
}

一、整体架构总览

这是昇腾CANN 算子双版本模板代码，一套代码兼容两种运行模式：

二、关键名词通俗解释

三、ACL 标准执行流程（固定顺序，不能乱）

ACL 模式是昇腾应用开发标准流程，创建 & 释放必须逆序：

四、两大分支使用场景区分

• 开发调试阶段：加编译宏 -DCCE_KT_TEST，用CCE仿真，不用插昇腾卡，快速调试内核逻辑。
• 上线部署阶段：不加宏，走ACL硬件模式，跑在真实昇腾卡上，发挥硬件加速能力。
  ai生成