AI Core编译基本用法【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit当开发者完成一个Kernel核函数的编写并在Host侧通过方式完成核函数的启动调用后就需要进行算子源码编译与算子运行接下来我们介绍通过bisheng编译器完成算子编译的基本用法。AI Core SIMD编译使用bisheng编译Ascend C源文件的基本命令如下其中--npu-archdav-npu architecture用于指定AI处理器架构版本bisheng source_file.asc -o output_file --npu-archdav-npu architecturebisheng可以通过-I path指定头文件路径、-L path 指定链接库路径、-l library指定需要链接的动态库或静态库、-Dmacrovalue指定宏定义示例如下bisheng add_kernel.asc -I path_to_include -L path_to_library -l library -o output_file -Dmacrovalue --npu-archdav-npu architectureAI Core SIMD的基本编译流程如下Host代码使用Host编译器编译成Host二进制AI Core SIMD代码分为Cube代码和Vector代码需分别编译成Cube二进制和Vector二进制先将Cube二进制和Vector二进制链接成Fatbin文件再与Host二进制合并生成可执行二进制。异构编译完整样例请参考LINK。// ----- add_kernel.asc ----- // 指定核函数在AI Core上执行 __global__ __vector__ add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) bisheng -c add_kernel.asc -o add_kernel.o --npu-archdav-xxxx bisheng -c main.cpp -o main.o -I${INSTALL_DIR}/include bisheng add_kernel.o main.o -o main // 或 bisheng main.cpp add_kernel.asc -o main --npu-archdav-xxxx[!NOTE] 说明${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。以root用户安装为例安装后文件默认存储路径为/usr/local/Ascend/cann。单独编译完整样例请参考LINK。bisheng异构编译默认采用的是全程序编译模式该模式要求单个源文件X.asc中编译的设备程序没有任何未解析的外部设备函数和变量引用。设备函数可以调用其他编译单元中定义的设备函数或访问其中定义的设备变量但必须在bisheng命令行中指定-dc编译选项才能启用不同编译单元间设备代码的链接功能。这种支持跨编译单元链接设备代码与符号的能力被称为单独编译。单独编译能实现更灵活的代码组织、缩短编译耗时并生成体积更小的可执行文件相较于全程序编译它会增加一定的编译构建复杂度。此外设备代码链接可能会影响程序性能这也是该模式未被设为默认模式的原因。链接时优化LTO可有效降低单独编译带来的性能损耗。单独编译的强制要求在一个编译单元中定义的非常量设备变量在其他编译单元中引用时必须使用extern关键字声明所有常量设备变量定义和跨单元引用时都必须使用extern关键字所有Ascend C源文件.asc都必须使用-dc选项编译。下方示例中add_compute.asc定义了变量和函数add_kernel.asc对其进行引用两个文件将分别编译最终链接为完整可执行文件。// ------- add_compute.asc ------- __gm__ int dev_var 5; __aicore__ void add_compute(...); // ------- add_kernel.asc ------- extern __gm__ int dev_var; __aicore__ void add_compute(...); __global__ __vector__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { dev_var 0; add_compute(...); }bisheng -dc add_compute.asc -o add_compute.o --npu-archdav-xxxx bisheng -dc add_kernel.asc -o add_kernel.o --npu-archdav-xxxx bisheng -c main.cpp -o main.o -I${INSTALL_DIR}/include bisheng add_compute.o add_kernel.o main.o -o program[!NOTE] 说明${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。以root用户安装为例安装后文件默认存储路径为/usr/local/Ascend/cann。动态库编译完整样例请参考LINK。bisheng -shared add_kernel.asc -o libadd_kernel.so -fPIC --npu-archdav-xxxx静态库编译完整样例请参考LINK。bisheng -lib add_kernel.asc -o libadd_kernel.a --npu-archdav-xxxxAI Core SIMT编译使用bisheng编译器编译Ascend C SIMT源文件时需要在AI Core SIMD编译的基本命令上增加--enable-simt示例如下bisheng source_file.asc -o output_file --npu-archdav-npu architecture --enable-simtAI Core SIMT的基本编译流程如下Host代码使用Host编译器编译成Host二进制AI Core SIMT代码编译成SIMT二进制先将SIMT二进制链接成Fatbin文件再与Host二进制合并生成可执行二进制。基本编译命令汇总|编译方式|AI Core SIMD编译命令| AI Core SIMT编译命令| |-|-|-| |异构编译| bisheng source_file.asc -o output_file --npu-archdav-npu architecture | bisheng source_file.asc -o output_file --npu-archdav-npu architecture--enable-simt| |独立编译|bisheng-dcsource_file.asc -o output_file.o--npu-archdav-npu architecture| bisheng-dcsource_file.asc -o output_file.o--npu-archdav-npu architecture--enable-simt| |动态库编译|bisheng-sharedsource_file.asc -o output_file.so--npu-archdav-npu architecture-fPIC| bisheng-sharedsource_file.asc -o output_file.so--npu-archdav-npu architecture--enable-simt-fPIC| |静态库编译|bisheng-libsource_file.asc -o output_file.a--npu-archdav-npu architecture| bisheng-libsource_file.asc -o output_file.a--npu-archdav-npu architecture--enable-simt|常用的编译选项选项是否必需说明-help否查看帮助。--npu-arch是编译时指定的AI处理器架构取值为dav-arch-version其中arch-version为NPU架构版本号各产品型号对应的架构版本号请通过对应关系表进行查询。--npu-soc否编译时指定的AI处理器型号npu-soc和npu-arch同时配置时优先采用npu-arch的配置。AI处理器的型号请通过如下方式获取针对如下产品在安装AI处理器的服务器执行npu-smi info命令进行查询获取Name信息。实际配置值为AscendName例如Name取值为xxxyy实际配置值为Ascendxxxyy。Atlas A2 训练系列产品 / Atlas A2 推理系列产品Atlas 200I/500 A2 推理产品Atlas 推理系列产品Atlas 训练系列产品针对如下产品在安装AI处理器的服务器执行npu-smi info -t board -iid-cchip_id命令进行查询获取Chip Name和NPU Name信息实际配置值为Chip Name_NPU Name。例如Chip Name取值为AscendxxxNPU Name取值为1234实际配置值为Ascendxxx_1234。其中id设备id通过npu-smi info -l命令查出的NPU ID即为设备id。chip_id芯片id通过npu-smi info -m命令查出的Chip ID即为芯片id。Ascend 950PR/Ascend 950DTAtlas A3 训练系列产品 / Atlas A3 推理系列产品-x否指定编译语言,如-x asc,表示指定为Ascend C编程语言。-o file否指定输出文件的名称和位置。-c否编译生成目标文件。-dc否编译生成relocatable目标文件。-shared--shared否编译生成动态链接库。-lib--cce-build-static-lib否编译生成静态链接库。编译器会将Device侧的代码进行编译链接生成Device侧二进制文件随后将该文件作为Host侧编译的输入进行编译最后链接生成静态链接库。-g否编译时增加调试信息。--sanitizer否编译时增加代码正确性校验信息。使用sanitizer选项时需要同步添加-g选项且不能在-O0场景下使用。注意启用该选项后GlobalTensor默认使用L2 Cache无法通过AscendC::SetL2CacheHint接口设置不使用L2 Cache的模式。-fPIC否告知编译器产生位置无关代码。-O否用于指定编译器的优化级别当前支持-O3-O2-O0。--run-modesim否sim模式链接时用户添加仿真模式对应的实现库实现代码在仿真模式下运行可以查看仿真相关日志方便用户性能调试。--enable-simt否SIMT编程场景指定SIMT方式编译。--cce-disable-asc-reserved-ubuf否禁用Ascend C接口使用预留UB空间。开启后依赖预留UB空间的Ascend C接口在对应芯片架构下不可用使用时编译报错。使用预留UB空间的API列表参考 使用预留UB空间的API。由于同一API使用预留UB空间的情况在不同NPU架构下有差异开启该编译选项后需要手动调整API调用方式或替换为不依赖预留UB空间的实现才能完成兼容性迁移。--cce-disable-vf-stack-reserved-ubuf否禁用SIMD VF栈预留的UB空间。开启后编译器不再预留该部分UB空间该空间可作为普通UB空间使用。针对 NPU架构版本2201此编译选项无实际效果针对 NPU架构版本3510此编译选项生效当用户使用此编译选项后编译器将无法使用预留的UB空间进行寄存器溢出的缓存需要用户保证寄存器不溢出。更多的编译命令和用法可以参见《毕昇编译器用户指南》。CMake方式编译项目中可以使用CMake来更简便地使用毕昇编译器编译Ascend C算子生成可执行文件、动态库、静态库或二进制文件。以下是CMake脚本的示例及其核心步骤说明# 1、设置编译配置变量 # CMAKE_ASC_ARCHITECTURES指定NPU架构版本当前默认值为dav-2201可通过-DCMAKE_ASC_ARCHITECTURESdav-xxxx覆盖默认值 set(CMAKE_ASC_ARCHITECTURES dav-2201 CACHE STRING NPU architecture) # 2、find_package(ASC)是CMake中用于查找和配置Ascend C编译工具链的命令 find_package(ASC) # 3、指定项目支持的语言包括ASCASC表示支持使用毕昇编译器对Ascend C编程语言进行编译 project(kernel_samples LANGUAGES ASC) # 4、使用CMake接口编译可执行文件、动态库、静态库、二进制文件 add_executable(demo add_custom.asc )编译AI Core SIMT代码时需要在编译选项中增加--enable-simt或通过CMAKE_ASC_ENABLE_SIMT设置为ON启用SIMT编译模式示例如下完整样例请参考LINK。# CMAKE_ASC_ARCHITECTURES指定NPU架构版本可通过-DCMAKE_ASC_ARCHITECTURESdav-xxxx覆盖默认值 set(CMAKE_ASC_ARCHITECTURES dav-3510 CACHE STRING NPU architecture) # CMAKE_ASC_ENABLE_SIMT启用SIMT编译模式需要在find_package(ASC)前设置 set(CMAKE_ASC_ENABLE_SIMT ON) find_package(ASC) project(kernel_samples LANGUAGES ASC) add_executable(demo add_custom.asc )以下是动态库、静态库编译示例同时展示如何将源文件切换为用语言ASC编译编译.asc文件生成动态库# 3、使用CMake接口编译动态库 add_library(kernel_shared SHARED cube.asc vector.asc mix.asc )编译.asc文件生成静态库# 3、使用CMake接口编译静态库 add_library(kernel_lib STATIC cube.asc vector.asc mix.asc )表1常用的CMAKE配置变量说明| 变量名称 | 配置说明 | |--|--| | CMAKE_BUILD_TYPE | 编译模式选项可配置为Release或Debug。Release版本不包含调试信息编译最终发布的版本。Debug版本包含调试信息便于开发者开发和调试。配置其他值时CMake会发出警告但不中止构建。 | | CMAKE_ASC_COMPILER | 指定Ascend C的编译器路径默认由CMake自动探测。路径不存在时配置阶段报错终止。 | | CMAKE_ASC_SOURCE_FILE_EXTENSIONS | ASC源文件扩展名列表默认为asc。 | | CMAKE_ASC_STANDARD | 指定编译使用的C标准版本默认值为17。设置低于17的值时CMake会发出警告。 | | CMAKE_ASC_ARCHITECTURES | 指定目标NPU架构版本取值为dav-arch-version如dav-2201、dav-3510。NPU运行模式下CMake会自动将该值以--npu-archvalue注入编译命令。 | | CMAKE_ASC_RUN_MODE | 指定算子运行模式支持npu默认在NPU设备上运行、cpuCPU调试模式、sim仿真模式。不同模式下CMake会自动调整链接库和编译选项。 | | CMAKE_ASC_FLAGS | 为所有构建类型DebugRelease等添加的通用编译选项。 | | CMAKE_ASC_FLAGS_DEBUG | 专用于Debug构建的编译选项默认为-O0 -g。 | | CMAKE_ASC_FLAGS_RELEASE | 专用于Release构建的编译选项默认为-O3 -DNDEBUG。 | | CMAKE_ASC_COMPILER_LAUNCHER | 编译命令的前置启动器如ccache、distcc未设置时直接调用编译器。 | | CMAKE_ASC_LINKER_LAUNCHER | 链接命令的前置启动器未设置时直接执行链接。 | | CMAKE_ASC_COMPILER_AR | 静态库归档工具默认使用系统ar工具。设置后静态库的归档命令将使用该工具替代默认ar。 | | CMAKE_ASC_COMPILER_LINKER | 链接驱动默认使用bisheng编译器作为链接驱动仅在有明确替换需求时设置。 | | CMAKE_ASC_ENABLE_SIMT | 是否启用SIMT编译模式设置为ON时在编译命令中自动注入--enable-simt选项默认为OFF。 | | CMAKE_INSTALL_PREFIX | 为CMake的内置变量用于指定CMake执行install命令时安装的路径前缀。|[!CAUTION]注意 在CMake工程中影响Ascend C语言初始化、编译器探测和编译规则生成的变量 需要在find_package(ASC)前设置或在执行cmake配置命令时通过-D指定。 这些变量在Ascend C语言初始化过程中会被读取若在find_package(ASC)之后再设置 可能不会影响已经生成的编译规则。其他编译相关说明内置编译宏开关内置编译宏开关列表如下ASCENDC_DUMP用于控制Dump开关默认开关打开开发者调用printf/DumpTensor/assert后会有信息打印需要注意直调工程的kernel文件内存在host函数如果在host函数内调用了printf接口也会触发kernel内的printf相关初始化动作进而影响kernel的执行性能)设置为0后表示开关关闭。示例如下# 关闭所有算子的printf打印功能 bisheng source_file.asc -o output_file --npu-archdav-npu architecture -DASCENDC_DUMP0ASCENDC_DEBUG用于控制Ascend C API的调测开关默认开关关闭增加该编译宏后表示开关打开此时接口内部的assert校验生效校验不通过会有assert日志打屏。开启该功能会对算子实际运行的性能带来一定影响通常在调测阶段使用。示例如下bisheng source_file.asc -o output_file --npu-archdav-npu architecture -DASCENDC_DEBUG当前ASCENDC_DEBUG功能支持的产品型号为 Atlas 推理系列产品 Atlas A2 训练系列产品/Atlas A2 推理系列产品ENABLE_CV_COMM_VIA_SSBUF用于控制是否使用SSBuffer以及UB到L1 Buffer的硬通道在涉及CV通信AIC和AIV或使用数据搬运API时需关注此选项。开启该选项可以提高相关API的性能或拓展使用更多功能。默认开关关闭设置为true后表示开关打开。示例如下bisheng source_file.asc -o output_file --npu-archdav-npu architecture -DENABLE_CV_COMM_VIA_SSBUFtrue仅在Ascend 950PR/Ascend 950DT支持该选项。从其它硬件平台移植到此平台的算子开关默认关闭以保持兼容性。在该平台新开发的算子以下场景需要打开使用矩阵计算Matmul高阶API且使用SetTensorScaleA等接口这些接口属于Ascend 950PR/Ascend 950DT新增的功能其内部实现使用了SSBuffer使用DataCopy接口从UB拷贝数据到L1 Buffer。NO_OVERLAP_IN_MULTI_REPEAT该编译选项用于在没有地址重叠的情况下移除不必要的内存同步指令以提升性能。针对Ascend 950PR/Ascend 950DT使用基础API的高维切分计算API时默认会插入内存同步指令以确保在地址重叠等复杂场景下的数据正确性但这些同步指令会带来性能开销。在追求极致性能的场景下如果您可以确定代码在任何情况下都不会发生内存重叠可以使用此选项。内置链接库毕昇编译器默认链接的库文件列表如下名称作用描述libascendc_runtime.aAscend C算子参数等组装库。libruntime.soRuntime运行库。libprofapi.soAscend C算子运行性能数据采集库。libunified_dlog.soCANN日志收集库。libmmpa.soCANN系统接口库。libascend_dump.soCANN维测信息库。libc_sec.soCANN安全函数库。liberror_manager.soCANN错误信息管理库。libascendcl.soacl相关接口库。高阶API常用链接库在使用高阶API时必须链接以下库因为这些库是高阶API功能所依赖的。在其他场景下可以根据具体需求选择是否链接这些库。 |链接库名称|作用描述|使用场景|动态库路径| |--|--|--|--| |libtiling_api.a|Tiling函数相关库。|使用高阶API相关的Tiling接口时需要链接。|${ASCEND_HOME_PATH}/lib64| |libregister.so|Tiling函数相关库。|使用高阶API相关的Tiling接口时需要链接。|${ASCEND_HOME_PATH}/lib64| |libgraph_base.so|基础数据结构和接口库。|调用ge::Shapege::DataType等基础结构体时需要链接。|${ASCEND_HOME_PATH}/lib64| |libplatform.so|硬件平台信息库。|使用PlatformAscendC相关硬件平台信息接口时需要链接。|${ASCEND_HOME_PATH}/lib64|【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考