【算子】1. Ascend C 最简工程实践记录实验代码目标创建一个最简 Ascend C 算子工程——Vector Add——从零编写、编译、在真实 NPU 上运行并验证结果。同时通过运行时 API 查询硬件平台参数理解代码与硬件的对应关系。环境搭建直接拉的CANN的镜像编译要给特权asc-devkit代码仓与cann 9.1.0有不兼容的修改后可完整编译学习材料获取https://gitcode.com/cann/asc-devkit什么是Ascend C-CANN社区版9.1.0-beta.1-昇腾社区Ascend C API列表-CANN商用版9.0.0-昇腾社区最小工程结构只需要两个文件simple/ ├── add.asc # kernel 实现 host 侧调用单文件合一 └── CMakeLists.txt # 编译配置add.asc 解析文件结构整个.asc文件由上到下分为三部分头文件区引入 kernel API、ACL 运行时、平台信息查询kernel 函数运行在 AI Core 上的设备端代码host 函数运行在 CPU 上的主控代码main 辅助函数头文件依赖关系#includekernel_operator.h// Ascend C kernel 侧 APIAdd, DataCopy, PipeBarrier...#includeacl/acl.h// ACL 运行时aclrtMalloc, aclrtMemcpy, aclrtSynchronizeDevice...#includetiling/platform/platform_ascendc.h// 平台信息查询PlatformAscendCManager, GetCoreMemSize...常见错误只引kernel_operator.h而忘记acl/acl.h会导致aclrtMallocHost、ACL_MEMCPY_HOST_TO_DEVICE等未定义。kernel_operator.h仅提供 kernel 侧 APIhost 侧的 ACL 运行时 API 需要单独引入。kernel 函数templateuint32_tblockLength__vector__ __global__voidadd_custom(__gm__float*x,__gm__float*y,__gm__float*z)template uint32_t blockLength—非类型模板参数NTTPC 标准语法。必须用模板而非普通参数因为Allocfloat, blockLength()的第二个参数需要编译期常量UB 静态分配。__vector__— 编译器内置关键字标记此 kernel 运行在 VectorAIV单元上。__global__— 编译器内置关键字标记此函数为 kernel 函数可从 host 调用。__gm__— 地址空间限定符标记指针指向 Global MemoryHBM。block_idx— 编译器内置变量当前 block 的索引0 ~ numBlocks-1。以上关键字均非 C 标准是 Bisheng 编译器在--asc-aicore-lang模式下识别的 Ascend C 语言扩展类似 CUDA 的__global__、__shared__等。kernel 内部的数据流Global Memory (HBM) │ │ GlobalTensor SetGlobalBuffer ← 将裸指针装箱绑定地址长度 │ DataCopy(UB←GM) ← MTE 搬运单元 │ PipeBarrierPIPE_ALL() ← 同步确保搬运完成 ▼ Unified Buffer (片上 SRAM, 191KB) │ │ Add(zLocal, xLocal, yLocal) ← Vector 计算单元 │ PipeBarrierPIPE_ALL() ← 同步确保计算完成 ▼ Unified Buffer │ │ DataCopy(GM←UB) ← MTE 搬运回写 │ PipeBarrierPIPE_ALL() ▼ Global Memory (HBM)关键约束计算 API 只能操作 Local MemoryUB不能直接操作 Global Memory。关键 API 详解SetGlobalBuffer — 指针装盒xGm.SetGlobalBuffer(xblock_idx*blockLength,blockLength);// ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ ^^^^^^^^^^^// Global Memory 起始地址裸指针 元素个数kernel 参数__gm__ float* x是裸指针没有边界信息。DataCopy等 API 需要知道数据长度才能正确搬运以及做边界检查。SetGlobalBuffer的作用就是将裸指针 长度封装为GlobalTensorT对象使其成为 Ascend C API 可安全操作的类型。它是GlobalTensor的成员函数定义在kernel_operator.h中。static_cast — 为什么能用而不需引头文件x[i]static_castfloat(i);// uint32_t → floatstatic_cast不是函数也不是宏是C 标准关键字。和for、if、return一样由编译器内置识别不需要任何#include就能使用。四种 C caststatic_cast、dynamic_cast、const_cast、reinterpret_cast全部是关键字不是头文件定义的。kernel 直调 vs 框架调用当前代码的模式叫kernel 直调Kernel Launch核心特征// 声明直接写在 .asc 源码中__global__voidadd_custom(__gm__float*x,...){...}// 调用用 语法就地发起add_custom20488,0(dx,dy,dz);三个标志特征本代码体现kernel 函数声明在源码中__global__ void add_custom(...)host 端用直接调用8, 0不做算子注册 / 不走框架没有ACLNN_REGISTER、没有模型文件与框架调用的对比直调模式当前 main() → add_custom8,0(...) → NPU 执行 └─ 调用方和 kernel 在同一个 .asc 源文件中 框架调用Aclnn 模式 main() → aclnnAddCustom(handle, ...) → 框架调度 → NPU 执行 └─ 需要单独编译 注册 op_type 通过 aclopCompile 或 ACLNN 接口调用直调适合快速验证和调试算子逻辑不需要算子注册、不用写 JSON 配置一个.asc文件包含全部。框架调用适合生产环境支持多算子组合、图编译优化。语法声明位置numBlocks, l2Ctrl不在任何头文件中声明。它是 Bisheng 编译器在--asc-aicore-lang模式下识别的内置语法编译器方言。类似地关键字定义来源作用__global__编译器内置标记 kernel 函数可从 host 调用__vector__编译器内置标记 kernel 运行在 Vector 单元SIMD__gm__编译器内置地址空间限定符指向 Global MemoryN, L2编译器内置kernel 启动语法block_idx编译器内置当前 block 索引在 Ascend C 头文件中grep不到任何#define __global__或#define __vector__——它们不是宏是编译器 parser 直接认识的关键字。这套语法设计借鉴了 CUDA 的grid, block, shared, stream。host 侧调用add_customblockLengthnumBlocks,0(xDevice,yDevice,zDevice);// 模板参数 Ascend C 扩展语法 运行时参数numBlocks, 0两个参数numBlocks 8启动 8 个 blockl2Ctrl 0L2 cache 控制标志默认行为numBlocks可以超过物理 AI Core 数量实测 500 个 block 在 40 核上正常运行runtime 调度器会自动排队——多余的 block 等待前面完成后补充执行。平台信息查询usingnamespaceplatform_ascendc;PlatformAscendC*platPlatformAscendCManager::GetInstance();plat-GetSocVersion();// 芯片型号枚举plat-GetCurNpuArch();// NPU 架构版本2201/3510plat-GetCoreNum();// 总 AI Core 数plat-GetCoreNumAic();// Cube 核心数plat-GetCoreNumAiv();// Vector 核心数plat-GetCoreMemSize(CoreMemType::UB,size);// 各存储单元容量该类定义在tiling/platform/platform_ascendc.h需要链接tiling_api和platform两个动态库。代码中的硬件关联点整个代码中与硬件直接绑定的一共 4 处其余都是标准 C1.__gm__— 地址空间限定符voidadd_custom(__gm__float*x,...)// kernel 函数参数标记指针指向 Global MemoryHBM/DDR编译器据此生成不同的访存指令。也是编译器内置关键字不是头文件定义的。2.LocalMemAllocatorHardware::UB— 指定使用 Unified BufferAscendC::LocalMemAllocatorAscendC::Hardware::UBubAllocator;Hardware::UB是一个枚举值声明在 AI Core 片上的 Unified Buffer 中分配内存。不同架构2201 / 3510的 UB 大小和 bank 结构不同编译器根据--npu-arch生成对应的内存布局。Allocfloat, 2048()在编译期从 UB 中预留 2048×48KB这部分占用会在编译时计入 UB 总分配量。3.numBlocks, 0— 指定用多少核add_customblockLengthnumBlocks,0(...);numBlocks8表示把数据切成 8 块最多并行运行在 8 个 AI Core 上。实际并行度受物理核数限制超额部分由 runtime 排队调度。4.--npu-archdav-2201— 编译目标架构# CMakeLists.txt set(CMAKE_ASC_ARCHITECTURES dav-2201 ...) target_compile_options(... --npu-archdav-2201)告诉 bisheng 为哪个硬件架构生成二进制。dav-2201和dav-3510的指令集、buffer 大小、bank 分组都不同不同架构的二进制互不兼容。完整映射关系CMakeLists.txt: --npu-archdav-2201 │ ▼ add.asc: __gm__ → 编译器知道是 HBM 地址空间 add.asc: Hardware::UB → 编译器知道用 256KB 片上 SRAM191KB 可用 add.asc: Allocfloat, N → 编译期从 UB 预留 N×4 字节 add.asc: DataCopy → 生成 MTE 搬运指令 add.asc: Add → 生成 Vector 计算单元指令 add.asc: PipeBarrier → 生成硬件同步信号 add.asc: 8, 0 → 指定一次启动 8 个 AI Core blockCMakeLists.txt 解析cmake_minimum_required(VERSION 3.16) # [通用] cmake 最低版本 set(CMAKE_ASC_ARCHITECTURES dav-2201 ...) # [硬件] 目标 NPU 架构 find_package(ASC REQUIRED) # [CANN] 注册 .asc 语言 加载编译器/头文件 project(add_project LANGUAGES ASC CXX) # [半通用] ASC 由 CANN 提供CXX 为标准 add_executable(add_demo add.asc) # [通用] 声明可执行目标 target_compile_options(add_demo PRIVATE # [硬件] 仅对 .asc 文件追加 --npu-arch $$COMPILE_LANGUAGE:ASC:--npu-arch${CMAKE_ASC_ARCHITECTURES} ) target_link_libraries(add_demo PRIVATE # [CANN] 平台查询所需的动态库 tiling_api platform )find_package(ASC)背后发生了什么1. 检查 $ASCEND_HOME_PATH 环境变量由 source set_env.sh 设置 2. 加载 config.cmake → 设置编译器路径、CANN 包路径、支持的 SOC 列表 3. 加载 func.cmake → 定义 npu_op_xxx 等高级函数算子工程用 4. 加载 intf.cmake → 添加 include 路径和链接路径 5. 注册 ASC 语言 → .asc 后缀 → bisheng 编译器 → CMakeASCInformation.cmake实际编译命令通过make VERBOSE1可见# 编译 .asc → .obisheng-fPIC--npu-archdav-2201-oadd.asc.o-c--asc-aicore-lang add.asc# ^^^^ ^^^^ ^^^^^^^^^^^^^^^^# 位置无关代码 输出 启用 Ascend C 方言# 链接 .o → 可执行文件bisheng add.asc.o-oadd_demo-ltiling_api-lplatform编译与运行# 1. 进入容器如果不在容器内dockerexec-itXXXXbash# 2. 加载环境每次新终端都要做source/usr/local/Ascend/cann/set_env.sh# 3. 编译cd/home/XXXX/ascend_C/demo/simplerm-rfbuildmkdirbuildcdbuild cmake-DCMAKE_ASC_ARCHITECTURESdav-2201..make-j# 4. 运行./add_demo# 5. 查看实际编译参数排查误编译makeVERBOSE1实际硬件查询结果在 XXXXXdav-XXXX上通过GetCoreMemSize运行时查询存储单元容量L0_A Buffer64 KBL0_B Buffer64 KBL0_C Buffer128 KBL1 Buffer511 KBL2 Cache192 MBUB (Unified Buffer)191 KB (196352 B)HBM (Global Memory)64 GBFixpipe Buffer2 KBBiasTable Buffer1 KBUB 返回 191KB 而非文档中的 256KB因为 API 返回的是 kernel 可用的动态分配空间已扣除系统预留8KB和最小 DCache32KB等固定开销。需要注意的点编译问题必须在容器内编译宿主机未安装CANNfind_package(ASC)找不到 ASCConfig.cmake。容器内有完整的 CANN 环境cmake 能正确找到。每次新终端必须source set_env.sh否则$ASCEND_HOME_PATH未设置cmake 报FATAL_ERROR。缺少 ACL 头文件是新手最容易犯的错误kernel_operator.h只提供 kernel 侧 API。aclrtMallocHost、ACL_MEMCPY_HOST_TO_DEVICE等定义在acl/acl.h中必须显式 include。aclInit/aclFinalize不能省略ACL 运行时需要正确初始化和释放否则内存泄漏或后续运行报错。编程问题kernel 内blockLength必须是编译期常量所以用模板参数传递不能用普通函数参数。numBlocks, 0中的 numBlocks 可以大于物理核数runtime 会自动调度排队无需手动限制。C 流输出不能直接打印枚举值GetCurNpuArch()返回NpuArch枚举需要用static_castuint32_t转换后才能std::cout。运行模式三种运行模式通过 cmake 参数切换NPU 真机默认无额外参数仿真-DCMAKE_ASC_RUN_MODEsim无需 NPU 硬件CPU 调试-DCMAKE_ASC_RUN_MODEcpu可在 CPU 上用 gdb编译参数验证始终用make VERBOSE1确认实际使用的编译器路径、--npu-arch值和链接库是否正确。命名空间platform_ascendc 中的类型需要命名空间前缀PlatformAscendC、CoreMemType等定义在platform_ascendc命名空间下要么写全称要么using namespace platform_ascendc;。