资讯动态

Ascend C SIMD 算子编译模式全解析:基本异构编译、单独编译与动态/静态库编译实战

发布时间:2026/10/14 1:47:23 来源:尧图企业网站定制
人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载导读本文基于 asc-devkit 仓库中examples/01_simd_cpp_api/02_features/04_compile下的四组编译样例系统讲解 Ascend C SIMD 算子的四种编译模式基本异构编译00_basic_compile、跨文件单独编译01_separate_compile、动态共享库编译02_dynamic_library_compile与静态库编译03_static_library_compile。读完本文你将掌握 bisheng 编译器的-c、-dc、-shared、-lib等核心选项的用法与适用场景理解.asc混合文件与纯 Host 侧.cpp的协作关系并能在 CMake 工程中通过CMAKE_ASC_ARCHITECTURES、CMAKE_ASC_RUN_MODE等变量一键完成 NPU 运行、CPU 调试与 NPU 仿真三种模式的编译切换。样例总览一个算子四种编译方式04_compile目录下的主 READMEREADME.md将四个样例组织为一张导航表它们均以Add 算子为载体输入/输出向量长度均为 16384单 block 处理长度 2048核心差异在于 Device 侧.asc代码的编译产物与组织方式样例目录编译模式核心编译选项Device 侧产物00_basic_compile基本异构编译-c.o目标文件与 Host 二进制合并为可执行文件01_separate_compile单独编译跨文件设备调用-dc多个.o目标文件跨编译单元链接02_dynamic_library_compile动态库编译-shared -fPIC.so动态共享库03_static_library_compile静态库编译-lib.a静态库四个样例支持的产品与 CANN 版本一致Ascend 950PR950DT 系列产品要求 CANN 9.1.0Atlas A3 与 Atlas A2 系列产品要求 CANN 9.0.0。前置知识.asc 混合文件与异构编译的基本流程Host Device 混合的.asc文件在 Ascend C SIMD 编程中.asc文件是Host Device 混合文件同一文件内既可以包含__vector__ __global__修饰的 Kernel 定义与内核启动调用Device 侧逻辑也可以包含供纯 Host 侧.cpp调用的普通函数。以基本编译样例的 add_kernel.asc 为例其核心结构如下template uint32_t blockLength __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { AscendC::InitSocState(); AscendC::GlobalTensorfloat xGm, yGm, zGm; xGm.SetGlobalBuffer(x block_idx * blockLength, blockLength); // ... 在 UB 上分配 LocalTensor经 DataCopy 从 GM 搬入 AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::SetFlagAscendC::HardEvent::MTE2_V(EVENT_ID0); AscendC::WaitFlagAscendC::HardEvent::MTE2_V(EVENT_ID0); AscendC::Add(zLocal, xLocal, yLocal, blockLength); // ... 结果经 DataCopy 写回 GM AscendC::PipeBarrierPIPE_ALL(); } void run_add_custom(float* x, float* y, float* z, uint32_t total_length, uint32_t num_blocks, aclrtStream stream) { add_customBLOCK_LENGTHnum_blocks, 0, stream(x, y, z); }该文件通过extern将内核启动函数run_add_custom暴露给纯 Host 侧的 main.cpp。Host 侧代码负责标准 ACL 流程aclInit→aclrtSetDevice→aclrtCreateStream→aclrtMalloc分配设备内存 →aclrtMemcpy搬运输入 → 调用内核 →aclrtSynchronizeStream等待完成 → 回拷结果 → 释放资源最后通过verify_result将算子输出与 Host 侧逐元素求和得到的 golden 值比对打印test pass!。异构编译的产物流根据 AI Core 编译基本用法 的说明AI Core SIMD 的编译流程为Host 代码使用 Host 编译器编译成 Host 二进制AI Core SIMD 代码由 Device 编译器编译成 SIMD 二进制先将 SIMD 二进制链接成 Fatbin 文件再与 Host 二进制合并生成可执行二进制。这决定了.asc与.cpp必须走不同的编译路径再通过链接阶段汇合也正是bisheng一条命令即可同时处理两类源文件的原因。四个样例的代码结构与调用关系可统一概括为纯 Host 的main.cpp通过extern声明→.asc混合文件暴露的启动函数 →内核启动 → Device 侧 Kernel 执行。基本编译与单独编译两种模式的差异如图 1、图 2 所示。图 1基本编译模式——单个.asc混合文件同时承载 Kernel 定义与启动调用图 2单独编译模式——Device 计算函数下沉到独立的add_compute.ascadd_kernel.asc通过extern跨文件调用模式一基本异构编译00_basic_compile基本异构编译是最常见的编译形态一个.asc混合文件 一个纯 Host.cpp分别编译成.o后链接为可执行二进制。完整说明见 00_basic_compile/README.md目录结构为├── 00_basic_compile │ ├── figures // 图示 │ ├── add_kernel.asc // 混合文件含 Kernel 定义、 内核调用 │ ├── CMakeLists.txt // 编译工程文件 │ ├── main.cpp // 纯 Host 侧代码通过 extern 调用 .asc 暴露的函数 │ └── README.md // 样例说明文档bisheng 命令行编译# 分步编译 bisheng -c add_kernel.asc -o add_kernel.o --npu-archdav-2201 bisheng -c main.cpp -o main.o -I${ASCEND_HOME_PATH}/include bisheng add_kernel.o main.o -o demo # 或一条编译命令 bisheng main.cpp add_kernel.asc -o demo --npu-archdav-2201 -I${ASCEND_HOME_PATH}/include关键点说明--npu-arch为必选参数指定 AI 处理器架构版本本样例支持dav-2201与dav-3510。各产品型号对应的架构版本号可在 simd_builtin_keywords.md 的对应关系表 中查询Ascend 950PR950DT 系列为 3510Atlas A3 系列为 2201Atlas A2 系列为 2201Atlas 200I/500 A2 推理产品为 3002Atlas 推理系列为 2002Atlas 训练系列为 1001。-I${ASCEND_HOME_PATH}/include提供acl/acl.h、kernel_operator.h等头文件的搜索路径ASCEND_HOME_PATH为 CANN 安装路径下的环境变量。分步编译时.asc与.cpp分别生成目标文件最后统一链接单命令方式则让 bisheng 自动识别文件类型、走完整个异构编译流程。CMake 方式编译样例的 CMakeLists.txt 展示了在 CMake 工程中启用 ASC 语言的标准步骤cmake_minimum_required(VERSION 3.16) # 1、设置编译配置变量指定 NPU 架构版本可通过 -DCMAKE_ASC_ARCHITECTURESdav-xxxx 覆盖 set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING NPU architecture: dav-2201, dav-3510) # 2、find_package(ASC) 用于查找并配置 Ascend C 编译工具链 find_package(ASC REQUIRED) # 3、指定项目支持的语言包括 ASC 和 CXX project(basic_compile LANGUAGES ASC CXX) # 4、使用 CMake 接口编译可执行文件.asc 与 .cpp 一并列入 add_executable(demo add_kernel.asc main.cpp) # 5、设置链接器语言为 ASC set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC) # 6、为 ASC 语言源文件注入 --npu-arch 编译选项 target_compile_options(demo PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-arch${CMAKE_ASC_ARCHITECTURES}) target_include_directories(demo PRIVATE $ENV{ASCEND_HOME_PATH}/include)其底层机制在工具链源码中可以得到印证find_package(ASC)会加载cmake/asc/asc_modules下的编译器探测与信息初始化脚本。CMakeDetermineASCCompiler.cmake 根据CMAKE_SYSTEM_PROCESSORx86_64 或 aarch64定位${ASCEND_HOME_PATH}/x86_64-linux/ccec_compiler/bin/或aarch64-linux目录下的bisheng编译器并将.asc注册为 ASC 语言的默认源文件扩展名、默认 C 标准设为 17。而 CMakeASCInformation.cmake 则会在检测到CMAKE_ASC_ARCHITECTURES时自动向全局 ASC 编译标志追加--npu-archvalue并将单文件编译规则拼装为bisheng ... -c --asc-aicore-lang source这正是 CMakeLists 中可以省略显式-c的原因。编译运行与三种运行模式在样例根目录执行source ${install_path}/cann/set_env.sh # 配置 CANN 环境变量${install_path} 默认为 /usr/local/Ascend mkdir -p build cd build cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # 默认 npu 模式 ./demo运行成功输出test pass!表示 Device 侧计算结果与 Host 侧 golden 值逐元素一致。CPU 调试与 NPU 仿真模式的编译cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..; make -j; # NPU 仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake否则旧的缓存变量会覆盖新的-D配置。两个核心编译选项的取值与含义如下表选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201对应 Atlas A2 与 Atlas A3 系列产品dav-3510对应 Ascend 950PR950DT 系列产品三种运行模式在 CMakeASCInformation.cmake 中的实现逻辑第 163–247 行值得深入理解模式合法性校验CMAKE_ASC_RUN_MODE必须是npu、cpu、sim三者之一否则配置阶段直接FATAL_ERROR。sim仿真模式仅支持dav-2002/dav-2201/dav-3510/dav-9201架构自动链接$ASCEND_HOME_PATH/tools/simulator/arch/lib下的runtime_camodel、npu_drv_camodel库并注入-Wl,-rpath与--disable-new-dtags。cpuCPU 调试模式按架构映射到对应的 CPU 模型 SoC 目录如dav-2201→Ascend910B1、dav-3510→Ascend950PR_9599追加-g -D_GLIBCXX_USE_CXX11_ABI0 --run-modecpu编译选项并链接ascendc_acl_stub、pem_davinci、cpudebug_cceprint、cpudebug、unified_dlog等一组调试/仿真库。模式二单独编译01_separate_compile当算子规模增大、希望将 Device 侧逻辑拆分到多个.asc文件时就进入了单独编译场景。完整说明见 01_separate_compile/README.md。为什么需要-dcbisheng 默认采用全程序编译模式该模式要求单个源文件X.asc中编译的设备程序没有任何未解析的外部设备函数和变量引用。当需要在多个.asc文件间进行跨文件设备代码调用时必须使用-dc选项启用单独编译模式才能支持不同编译单元间设备代码与符号的链接。单独编译能实现更灵活的代码组织、缩短编译耗时并生成体积更小的可执行文件代价是构建复杂度上升且设备代码链接可能带来一定性能影响——这正是它未被设为默认模式的原因链接时优化 LTO 可有效降低这部分损耗见 AI Core 编译基本用法。单独编译有三条强制要求在一个编译单元中定义的非常量设备变量在其他编译单元中引用时必须使用extern关键字声明所有常量设备变量定义和跨单元引用时都必须使用extern关键字所有 Ascend C 源文件.asc都必须使用-dc选项编译。跨文件调用的代码组织本样例把计算逻辑下沉到独立的 Device 侧执行函数文件 add_compute.asc__aicore__ void add_compute( AscendC::LocalTensorfloat zLocal, AscendC::LocalTensorfloat xLocal, AscendC::LocalTensorfloat yLocal, uint32_t length) { AscendC::Add(zLocal, xLocal, yLocal, length); }混合文件 add_kernel.asc 则通过extern声明该 Device 函数并在 Kernel 内调用extern __aicore__ void add_compute( AscendC::LocalTensorfloat zLocal, AscendC::LocalTensorfloat xLocal, AscendC::LocalTensorfloat yLocal, uint32_t length); // ... add_custom Kernel 内部 add_compute(zLocal, xLocal, yLocal, blockLength);由此形成三层结构main.cpp纯 Hostextern调用启动函数→add_kernel.asc混合文件Kernel 定义 启动→add_compute.asc纯 Device 计算函数经extern __aicore__暴露。三个文件分别编译后再链接。bisheng 命令行编译bisheng -dc add_compute.asc -o add_compute.o --npu-archdav-2201 bisheng -dc add_kernel.asc -o add_kernel.o --npu-archdav-2201 bisheng -c main.cpp -o main.o -I${ASCEND_HOME_PATH}/include bisheng add_compute.o add_kernel.o main.o -o demo注意所有.asc文件都必须加-dcHost 侧.cpp仍用-c。CMake 方式编译CMake 写法与基本编译几乎一致唯一区别是在 ASC 编译选项中加入-dc见 01_separate_compile/CMakeLists.txtadd_executable(demo add_compute.asc add_kernel.asc main.cpp) set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC) target_compile_options(demo PRIVATE $$COMPILE_LANGUAGE:ASC:-dc --npu-arch${CMAKE_ASC_ARCHITECTURES})编译运行、三种运行模式与执行结果test pass!均与模式一相同。模式三动态库编译02_dynamic_library_compile当希望将算子的 Device 侧实现封装为可独立分发、动态加载的组件时可选择动态库编译。完整说明见 02_dynamic_library_compile/README.md。其思路是.asc混合文件被编译为独立的动态共享库.soHost 侧.cpp通过链接该.so完成调用从而解耦算子实现与调用方。bisheng 命令行编译使用-shared选项并配合-fPIC生成位置无关代码bisheng -shared add_kernel.asc -o libadd_kernel.so -fPIC --npu-archdav-2201CMake 方式编译在 CMake 工程中只需将add_executable换成add_library(... SHARED ...)并将可执行目标与库目标链接见 02_dynamic_library_compile/CMakeLists.txt# 将 .asc 文件编译为动态共享库 add_library(add_kernel SHARED add_kernel.asc) set_target_properties(add_kernel PROPERTIES LINKER_LANGUAGE ASC) # Host 可执行程序链接该动态库 add_executable(demo main.cpp) set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC) target_compile_options(add_kernel PRIVATE $$COMPILE_LANGUAGE:ASC:--npu-arch${CMAKE_ASC_ARCHITECTURES} -fPIC) target_link_libraries(demo PRIVATE add_kernel)从工具链实现看CMakeASCInformation.cmake 为 ASC 语言配置了共享库创建规则CMAKE_ASC_CREATE_SHARED_LIBRARY走 bisheng 链接驱动与运行时-Wl,-rpath规则因此 CMake 工程无需手工拼装-shared链接参数-fPIC则通过CMAKE_ASC_COMPILE_OPTIONS_PIC统一管理。模式四静态库编译03_static_library_compile与动态库对应静态库编译将.asc编译为独立的静态库.a链接进最终可执行文件。完整说明见 03_static_library_compile/README.md。bisheng 命令行编译使用-lib选项等价于--cce-build-static-libbisheng -lib add_kernel.asc -o libadd_kernel.a --npu-archdav-2201从选项说明可知-lib的内部流程是编译器先将 Device 侧代码编译链接生成 Device 侧二进制文件再将该文件作为 Host 侧编译的输入进行编译最后链接生成静态链接库。CMake 方式编译# 将 .asc 文件编译为静态库 add_library(add_kernel STATIC add_kernel.asc) set_target_properties(add_kernel PROPERTIES LINKER_LANGUAGE ASC) add_executable(demo main.cpp) set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC) target_link_libraries(demo PRIVATE add_kernel)对应的 03_static_library_compile/CMakeLists.txt 还展示了target_include_directories指向$ENV{ASCEND_HOME_PATH}/include的 Host 头文件配置。工具链层面静态库归档由CMAKE_ASC_ARCHIVE_CREATE规则驱动默认使用系统ar可通过CMAKE_ASC_COMPILER_AR替换完成后以CMAKE_RANLIB收尾索引。四种模式的对比与选用建议维度基本异构编译单独编译动态库编译静态库编译核心选项-c-dc-shared -fPIC-lib跨文件设备调用不支持单编译单元支持extern __aicore__不支持不支持产物可执行文件可执行文件.so.a典型场景单文件小算子、快速验证算子代码模块化拆分、缩短编译耗时算子实现独立分发、动态加载算子实现打包进最终二进制、无运行时依赖构建复杂度最低中等需-dc extern 规范中等需处理链接路径低选用建议快速验证与单文件算子直接用基本异构编译一条 bisheng 命令或最小 CMake 工程即可Device 逻辑需要拆分复用、或单个文件编译过慢时切换到单独编译务必为所有.asc加上-dc并遵守 extern 声明规范常量设备变量定义与引用均需 extern需要将算子作为可插拔组件发布、热更新或供不同宿主程序复用时选择动态库编译追求部署简单、目标二进制自包含时选择静态库编译代价是产物体积增大、无法热更新。延伸bisheng 常用编译选项与 CMAKE 变量速查以下常用编译选项整理自 AI Core 编译基本用法选项是否必需说明--npu-arch是指定 AI 处理器架构取值为dav-arch-version产品与架构对应关系见 simd_builtin_keywords.md#npu-arch-c否编译生成目标文件-dc否编译生成 relocatable 目标文件单独编译模式-shared/--shared否编译生成动态链接库-lib/--cce-build-static-lib否编译生成静态链接库-fPIC否产生位置无关代码动态库必需-O否优化级别支持-O3、-O2、-O0-g否增加调试信息--enable-simt否指定 SIMT 方式编译SIMD 场景不需要--run-modesim否链接仿真实现库便于性能调试CMake 工程中影响 ASC 语言初始化的变量必须在find_package(ASC)之前设置或通过 cmake 配置命令的-D指定否则可能不影响已生成的编译规则这一点在 AI Core 编译基本用法 中有明确警告。常用变量包括CMAKE_ASC_ARCHITECTURES目标 NPU 架构取值为dav-arch-versionNPU 模式下自动注入--npu-archCMAKE_ASC_RUN_MODEnpu默认/cpu/sim不同模式自动调整链接库与编译选项CMAKE_ASC_STANDARDC 标准默认 17低于 17 时发出警告CMAKE_ASC_FLAGS/CMAKE_ASC_FLAGS_DEBUG默认-O0 -g/CMAKE_ASC_FLAGS_RELEASE默认-O3 -DNDEBUG编译选项CMAKE_ASC_COMPILER编译器路径默认自动探测路径不存在时配置阶段报错CMAKE_ASC_ENABLE_SIMT置 ON 时自动注入--enable-simtCMAKE_ASC_COMPILER_AR/CMAKE_ASC_COMPILER_LINKER静态库归档工具与链接驱动。若需在无 NPU 设备的开发环境中快速验证算子逻辑可参考 quick_start.md 的环境准备 完成 CANN 开发套件包安装与set_env.sh配置后直接使用cpu或sim模式编译运行本组样例。小结asc-devkit 的04_compile样例组以同一个 Add 算子演示了 Ascend C SIMD 四种编译模式的完整链路基本异构编译解决单文件快速编译运行-dc单独编译解决跨文件设备代码复用与模块化-shared与-lib则分别提供动态/静态库形态的封装分发。配合 CMake 的 ASC 语言工具链CMakeDetermineASCCompiler.cmake 与 CMakeASCInformation.cmake开发者只需维护少量配置变量即可在同一套工程内切换 NPU 运行、CPU 调试与 NPU 仿真三种模式是理解 Ascend C 算子工程化编译与集成的理想入门素材。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐Ascend C SIMT 算子四类编译模式实战基本异构编译、单独编译与动态/静态库编译全解析Ascend C SIMT 算子四类编译模式实战基本异构编译、单独编译与动态/静态库编译全解析 本文基于 CANN asc devkit 仓库中 SIMT A人工智能深度学习算子库CANNAscendAscend C SIMT 编译模式实战基本异构编译、单独编译与库编译基于 asc-devkit 编译样例Ascend C SIMT 编译模式实战基本异构编译、单独编译与库编译基于 asc devkit 编译样例 导读 本指南以 examples/03_sim人工智能深度学习算子库CANNAscenduni-app x 编译器全解析uts/uvue 编译链路、条件编译、编译缓存与静态资源处理实战指南uni app x 编译器全解析uts/uvue 编译链路、条件编译、编译缓存与静态资源处理实战指南 uni app x 的编译器由 uts 语言编译器与 u示例工程前端移动开发跨平台上一篇Ant Design Progress 组件设计指南从任务进度理解到圆形与微型进度条的应用解析下一篇Ant Design Progress 语义化 DOM 定制指南用 classNames 与 styles 精确控制进度条样式创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考

读完文章,也想定制专属网站?

尧图设计师 24 小时内与您沟通定制方案

免费获取报价 →
↑