企业官网建设流程全解析

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

导读

本文基于 asc-devkit 仓库中examples/01_simd_cpp_api/02_features/05_aclrtc目录下的样例，系统讲解Aclrtc（运行时编译）技术：在 Host 侧把字符串形式的 Ascend C 核函数源码交给 aclrtc 接口现场编译为 deviceELF，再通过 aclrt 运行时接口完成加载与启动。读完本文，你将掌握aclrtcCreateProg→aclrtcCompileProg→aclrtcGetBinData的编译链路、aclrtBinaryLoadFromData→aclrtLaunchKernelWithArgsArray的执行链路，以及模板核函数 mangled name 的获取方法，能够直接复现两个可运行的样例并理解其底层实现原理。

一、什么是 Aclrtc 运行时编译

传统自定义算子开发需要编写算子工程、预编译算子包（.run），再经算子注册流程才能被框架调用。而Aclrtc（runtime compilation）模式彻底改变了这一路径：核函数源码以字符串形式直接嵌入 Host 代码，程序运行时才在 Host 侧调用 aclrtc 接口完成编译，随后把编译产物加载到 NPU 并启动执行。这种"源码即参数、编译在运行时"的方式无需提前编译算子包，非常适合快速验证和原型开发。

从仓库源码看，aclrtc 的能力声明位于 tools/include/acl/acl_rt_compile.h，其实现位于 tools/aclrtc/acl_rt_compile.cpp。该实现通过dlopen动态加载 CANN 安装目录下毕昇编译器的libasrtc.so（路径为${cannPath}/tools/bisheng_compiler/lib/libasrtc.so），再以函数指针方式调用asrtcCreateProgram、asrtcCompileProgram、asrtcGetDeviceELF等接口——也就是说，aclrtc 是 Host 侧运行时与毕昇编译器之间的适配层，编译选项、源码字符串、产物二进制都在这一层完成传递。

一个值得注意的细节：aclrtc 会根据传入的程序名后缀判断编译类型（GetAclrtcTypeWithSuffix，见acl_rt_compile.cpp），.asc走 Ascend C 类型（对应枚举ACL_RTC_TYPE_ASC），.aicpu走 AI CPU 类型、.cce走 CCE 类型。目前源码注释明确aicpu compile is not supported yet，即 AI CPU 编译尚不支持。

二、样例总览

examples/01_simd_cpp_api/02_features/05_aclrtc/README_en.md列出了目录下的两个样例：

两个样例都需要CANN 9.1.0 及以上版本。它们覆盖了 aclrtc 的两类典型场景：普通（非模板）核函数的直接名字查找，以及模板核函数的aclrtcAddNameExpr+aclrtcGetLoweredName名字规约流程，下文分别详解。

三、样例一：rtc_hello_world——printf 打印与空参核函数启动

3.1 目录结构与核函数实现

该样例的文件组织如下（见 README_en.md）：

rtc_hello_world ├── CMakeLists.txt // 编译工程文件 ├── rtc_hello_world.cpp // Host 侧代码（含 RTC 核函数源码字符串） └── README.md // 样例说明文档

核函数源码以字符串形式嵌入 Host 代码。在 rtc_hello_world.cpp 中可以看到，它通过R""""(...)"""原始字符串字面量承载一段完整的核函数：

const char* src = R""""( #include "utils/debug/asc_printf.h" extern "C" __global__ __vector__ void hello_world() { printf("Hello World!!!\n"); } )"""";

其中__global__ __vector__是 Ascend C 的核函数声明关键字，printf输出结果，extern "C"保证导出名字不被 C++ 名字修饰（name mangling）改写，从而在执行阶段可以直接用"hello_world"查找函数句柄。

3.2 编译阶段：aclrtc 接口调用链

Host 侧核心流程分为编译阶段与执行阶段两个阶段。编译阶段的接口调用顺序如下：

1. aclrtcCreateProg—— 创建编译程序对象，传入核函数源码字符串。样例中第四个参数（程序名）传的是"hello_world.asc"，正是利用.asc后缀让底层识别为 Ascend C 编译类型；末尾的0, nullptr, nullptr分别对应numHeaders、headers、includeNames，按 acl_rt_compile.h 的注释，当前必须传 0 与 NULL，Header 支持尚未实现。
2. aclrtcCompileProg—— 执行运行时编译。通过 options 传入--npu-arch指定 NPU 架构：

const char* options[] = { "--npu-arch=" ACL_RTC_NPU_ARCH, // 默认 "dav-2201" }; aclError ret = aclrtcCompileProg(prog, 1, options);

3. aclrtcGetCompileLogSize/aclrtcGetCompileLog——可选，编译失败时获取错误日志。样例在编译返回值非成功时按"先取大小、再取内容"的方式打印日志：

if (ret != ACL_SUCCESS) { size_t size = 0; (void)aclrtcGetCompileLogSize(prog, &size); char log[size] = {0}; (void)aclrtcGetCompileLog(prog, log); printf("Compile Error Log : %s", log); }

4. aclrtcGetBinDataSize/aclrtcGetBinData—— 获取编译产物（deviceELF）的二进制大小与数据：

size_t binDataSizeRet; ASCENDC_CHECK(aclrtcGetBinDataSize(prog, &binDataSizeRet)); std::vector<char> deviceELF(binDataSizeRet); ASCENDC_CHECK(aclrtcGetBinData(prog, deviceELF.data()));

5. aclrtcDestroyProg—— 编译与运行全部结束后销毁编译程序对象。

3.3 执行阶段：aclrt 接口调用链

执行阶段先把设备初始化好（aclInit→aclrtSetDevice(0)→aclrtCreateStream），再走如下关键步骤：

• aclrtBinaryLoadFromData—— 将编译产物加载到设备。加载时需要指定二进制魔数类型：

aclrtBinaryLoadOptions loadOption; loadOption.numOpt = 1; aclrtBinaryLoadOption option; option.type = ACL_RT_BINARY_LOAD_OPT_MAGIC; option.value.magic = ACL_RT_BINARY_MAGIC_ELF_AICORE; loadOption.options = &option; ASCENDC_CHECK(aclrtBinaryLoadFromData(deviceELF.data(), binDataSizeRet, &loadOption, &binHandle));

• aclrtBinaryGetFunction—— 通过核函数名hello_world获取核函数句柄：

const char* funcName = "hello_world"; ASCENDC_CHECK(aclrtBinaryGetFunction(binHandle, funcName, &funcHandle));

• aclrtLaunchKernelWithArgsArray—— 本样例核函数无参数，传入空参数数组并启动核函数。注意核函数以 8 个 Block 并发执行：

uint32_t numBlocks = 8; void* kernelArgs[] = {}; ASCENDC_CHECK(aclrtLaunchKernelWithArgsArray(funcHandle, numBlocks, stream, nullptr, kernelArgs));

• aclrtSynchronizeStream—— 等待 stream 上核函数执行完成。

执行结束后依次aclrtBinaryUnLoad、aclrtDestroyStream、aclrtResetDevice、aclFinalize完成资源释放，最后aclrtcDestroyProg(&prog)销毁编译程序对象。

3.4 编译运行与执行结果

在样例根目录（即rtc_hello_world目录）执行：

source ${install_path}/cann/set_env.sh # 配置 CANN 环境变量 mkdir -p build && cd build; # 创建并进入 build 目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程 ./demo # 执行样例

${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。环境变量配置方式可参考 快速入门。

执行成功时输出如下，8 个 AIV Block 各打印一行：

[AIV Block 0/8] Hello World!!! [AIV Block 1/8] Hello World!!! [AIV Block 2/8] Hello World!!! [AIV Block 3/8] Hello World!!! [AIV Block 4/8] Hello World!!! [AIV Block 5/8] Hello World!!! [AIV Block 6/8] Hello World!!! [AIV Block 7/8] Hello World!!!

四、样例二：rtc_template_add——模板核函数 + 静态 Tensor 编程

4.1 样例规格

rtc_template_add/README_en.md 给出了完整的样例规格表：

计算公式为z = x + y。该样例的核心价值在于演示模板核函数在 aclrtc 下的处理方式：由于Kernel::add_custom<float>是模板实例化出来的符号，其最终符号名会被 C++ 编译器改写（mangling），无法像hello_world那样直接用源码名字查找，因此必须借助aclrtcAddNameExpr与aclrtcGetLoweredName两个接口完成"源码名 → mangled name"的映射。

4.2 Kernel 实现：LocalMemAllocator 与 SetFlag/WaitFlag

核函数源码同样以字符串嵌入 Host 代码，其实现采用静态 Tensor 编程范式（见 rtc_template_add.cpp）：

__aicore__ inline void Process() { AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator; AscendC::LocalTensor<float> xLocal = ubAllocator.Alloc<float, TOTAL_LENGTH>(); AscendC::LocalTensor<float> yLocal = ubAllocator.Alloc<float, TOTAL_LENGTH>(); AscendC::LocalTensor<float> zLocal = ubAllocator.Alloc<float, TOTAL_LENGTH>(); AscendC::DataCopy(xLocal, xGm, TOTAL_LENGTH); AscendC::DataCopy(yLocal, yGm, TOTAL_LENGTH); // MTE2→V 同步：确保 DataCopy 完成后再计算 AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); AscendC::Add(zLocal, xLocal, yLocal, TOTAL_LENGTH); // V→MTE3 同步：确保 Add 完成后再将结果拷出 AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); AscendC::DataCopy(zGm, zLocal, TOTAL_LENGTH); }

要点说明：

• LocalMemAllocator 静态内存管理：LocalMemAllocator<Hardware::UB>在 Unified Buffer 上按编译期常量TOTAL_LENGTH（1024）静态分配xLocal、yLocal、zLocal三块 LocalTensor，内存布局在编译期即可确定，避免运行时动态申请的开销。
• SetFlag/WaitFlag 事件同步：搬运单元 MTE2 写入完成后，通过SetFlag<HardEvent::MTE2_V>+WaitFlag<HardEvent::MTE2_V>确保DataCopy在计算单元 V 执行Add之前完成；计算完成后同样以V_MTE3事件对同步，保证结果在DataCopy(zGm, zLocal, ...)搬出前就绪。
• 核函数入口模板：Kernel::add_custom<T>是函数模板，Init中通过SetGlobalBuffer把三个GM_ADDR参数绑定为 GlobalTensor，Process完成"搬入 → 计算 → 搬出"全流程，末尾PipeBarrier<PIPE_ALL>()保证流水线整体同步。

4.3 模板核函数名的注册与获取

编译阶段与 hello_world 样例的差异集中在这里：

aclrtcProg prog; ASCENDC_CHECK(aclrtcCreateProg(&prog, src, "add_custom", 0, nullptr, nullptr)); const char* kernelNameExpr = "Kernel::add_custom<float>"; ASCENDC_CHECK(aclrtcAddNameExpr(prog, kernelNameExpr)); // 注册需要导出的模板核函数名 const char* options[] = { "--npu-arch=" ACLRTC_NPU_ARCH }; ASCENDC_CHECK(aclrtcCompileProg(prog, 1, options)); // ... aclrtcGetBinDataSize / aclrtcGetBinData ... const char* manglingName = ""; ASCENDC_CHECK(aclrtcGetLoweredName(prog, kernelNameExpr, &manglingName)); // 获取 mangled name

依据 acl_rt_compile.h 的接口契约：aclrtcAddNameExpr传入的 name expression 必须与后续aclrtcGetLoweredName传入的字符串完全一致，aclrtcGetLoweredName返回的 mangled name 内存在aclrtcDestroyProg销毁程序时释放。执行阶段随后用该 mangled name 查找核函数句柄：

ASCENDC_CHECK(aclrtBinaryGetFunction(binHandle, manglingName, &funcHandle));

4.4 数据准备与精度校验（纯 C++ 实现）

与 hello_world 样例不同，rtc_template_add 的核函数有 3 个参数（x、y、z的 Device 内存指针），因此执行阶段先通过aclrtMallocHost/aclrtMalloc（使用ACL_MEM_MALLOC_HUGE_FIRST分配策略）分别分配 Host 与 Device 内存，用aclrtMemcpy完成ACL_MEMCPY_HOST_TO_DEVICE的数据搬运，再构造参数数组启动：

void* kernelArgs[] = { static_cast<void*>(&xDevice), static_cast<void*>(&yDevice), static_cast<void*>(&zDevice), }; ASCENDC_CHECK(aclrtLaunchKernelWithArgsArray(funcHandle, numBlocks, stream, nullptr, kernelArgs));

数据生成与精度校验完全在 Host 侧 C++ 内完成，不依赖外部脚本。工具函数封装在 data_utils.h 中：

• GenTestData(x, y, golden, len)：生成[1.0, 100.0)区间的随机浮点输入，并同步计算黄金结果golden[i] = x[i] + y[i]；
• VerifyResult(output, golden, len)：按绝对容差 1e-5 + 相对容差 1e-4（ABSOLUTE_TOL、RELATIVE_TOL）逐元素比较，输出error ratio并统计误差率，误差率不超过ERROR_TOL（1e-4）即判定通过。

执行成功时输出：

test pass!

4.5 编译运行

在rtc_template_add样例根目录执行：

source ${install_path}/cann/set_env.sh mkdir -p build && cd build; cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 支持 dav-3510, dav-2201 ./demo

五、构建选项与运行约束

两个样例的 CMake 工程均基于ASCEND_HOME_PATH环境变量定位 CANN 头文件与动态库（见 rtc_hello_world/CMakeLists.txt 与 rtc_template_add/CMakeLists.txt）：

target_include_directories(demo PRIVATE $ENV{ASCEND_HOME_PATH}/include) target_link_directories(demo PRIVATE $ENV{ASCEND_HOME_PATH}/lib64) target_link_libraries(demo PRIVATE ascendcl acl_rtc)

即链接ascendcl与acl_rtc两个库。CMAKE_ASC_ARCHITECTURES通过target_compile_definitions以ACL_RTC_NPU_ARCH/ACLRTC_NPU_ARCH宏传入源码，最终拼进--npu-arch=编译选项。

关键构建选项说明：

⚠️ 注意：Aclrtc 仅支持 NPU 运行模式，不支持 CPU 调试或 sim 仿真模式。这一限制意味着调试核函数必须在真实 NPU 环境上完成，无法借助 CPU 模式或仿真器缩短验证周期。

六、底层实现原理：aclrtc 是如何工作的

结合 tools/aclrtc/acl_rt_compile.cpp，可以梳理出 aclrtc 的完整工作机理：

1. 动态加载编译器运行时。LoadExtraLib()通过dladdr反查libacl_rtc.so所在路径推导 CANN 安装目录，再dlopen毕昇编译器的libasrtc.so，并用dlsym解析出 9 个asrtc*符号（asrtcCreateProgram、asrtcCompileProgram、asrtcGetDeviceELF等）。因此运行时编译能力最终由 CANN 自带的毕昇编译器提供。

2. 编译选项的自动注入。aclrtcCompileProg在用户 options 之外会自动补充三类选项：语言类型选项-xasc（Ascend C 类型）、标准选项-std=c++17，以及 4 个头文件搜索路径（${cannPath}/include、${cannPath}/compiler/tikcpp/tikcfw及其interface、impl子目录）。这意味着字符串中的核函数源码可以直接#include "kernel_operator.h"等头文件，无需用户显式配置 include 路径。

3. 统一的错误码映射。底层asrtcResult通过ccecRet2AclrtcRet映射表转换为 aclrtc 的错误码（如ACL_ERROR_RTC_COMPILATION = 176003、ACL_ERROR_RTC_LINKING = 176004等，定义于acl_rt_compile.cpp顶部），上层应用只需判断返回值并配合aclrtcGetCompileLog获取详细日志即可定位问题。

4. 编译产物流转。aclrtcCompileProg成功后，aclrtcGetBinDataSize/aclrtcGetBinData取出的 deviceELF 二进制，直接交给执行阶段的aclrtBinaryLoadFromData加载——两条链路通过"内存中的二进制数据"而非磁盘文件衔接，这正是 aclrtc 与离线编译（生成 .o / .run 文件）模式的本质区别。

七、总结与适用场景

通过两个样例可以总结出 aclrtc 编程的核心套路：

1. 字符串即源码：把 Ascend C 核函数写成R""""(...)"""字符串嵌入 Host 代码；
2. 编译即运行：aclrtcCreateProg→（模板场景加aclrtcAddNameExpr）→aclrtcCompileProg→aclrtcGetBinDataSize/GetBinData，失败时用aclrtcGetCompileLogSize/GetCompileLog取日志；
3. 加载即执行：aclrtBinaryLoadFromData→aclrtBinaryGetFunction（模板场景用aclrtcGetLoweredName得到的 mangled name）→aclrtLaunchKernelWithArgsArray→aclrtSynchronizeStream；
4. 收尾清理：卸载二进制、销毁流/设备、aclrtcDestroyProg。

Aclrtc 尤其适合快速验证算子逻辑、原型开发、算法调优等无需打包部署算子的场景，但需注意其仅支持 NPU 运行模式、不支持 CPU 调试与仿真，且对模板核函数必须走AddNameExpr+GetLoweredName的名字规约流程。读者可以在此基础上，把 hello_world 的打印流程与 add_custom 的静态 Tensor 编程范式组合，构建属于自己的运行时编译算子验证框架。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend