异构编译
编译选项
常用的编译选项说明如下:
选项 |
是否必需 |
说明 |
---|---|---|
-help |
否 |
查看帮助 |
-o <file> |
否 |
Write output to <file> |
-O |
否 |
编译优化等级,-O2,-O0 |
-fPIC |
否 |
告知编译器产生与位置无关代码 |
--shared |
否 |
编译生成动态链接库 |
--cce-auto-sync |
否 |
开启自动同步 |
--std=c++17 |
否 |
结合Ascend C,需要使用C++17标准 |
-xcce |
否 |
代码中包含设备侧代码,但后缀名不是.cce时需添加 |
产品选项说明如下:
选项 |
是否必需 |
说明 |
---|---|---|
--cce-soc-version |
是 |
昇腾AI处理器的型号,配置后生成对应AI处理器型号的二进制。 如果无法确定具体的AI处理器型号,则在安装AI处理器的服务器执行npu-smi info命令进行查询,在查询到的“Name”前增加Ascend信息,例如“Name”对应取值为xxxyy,应配置为 --cce-soc-version=Ascendxxxyy。 |
--cce-soc-core-type |
是 |
生成对应核的二进制,参数取值为VecCore、CubeCore、AICore。
|
异构编译
毕昇编译器支持异构编译,当需要编写包含异构程序时,可按如下步骤进行:
- 从CANN发布包中获取毕昇编译器程序(bisheng)的路径,并设置环境变量。
# 设置编译器环境变量,如: $export PATH=${INSTALL_DIR}/compiler/ccec_compiler/bin:$PATH
- 代码开发。编写异构程序的基本流程如下:
- include必需的AscendCL Runtime(运行时管理)头文件。
- 通过Runtime API创建device和stream。
- 申请设备内存,按需将输入数据拷贝至设备内存。
- 编写设备kernel函数。
- 通过<<<>>>异构调用接口运行设备kernel函数。
- 按需将输出数据拷贝至主机内存,并释放设备内存。
- 获取runtime的动态链接库以及头文件的位置。在本示例中,编译依赖和所在位置如下:
编译依赖
位置
以下简称
runtime动态链接库
${INSTALL_DIR}/lib64/
$RT_LIB
runtime头文件
${INSTALL_DIR}/include
$RT_INC
- 编写并执行编译命令,一步编译生成可执行文件。下文场景示例为开发环境和运行环境一致的情况。
// 简化写法,无意义场景不可运行,仅演示编译使用方法 #include <stdio.h> #ifdef ASCENDC_CPU_DEBUG #define __aicore__ #else #define __aicore__ [aicore] #endif __global__ __aicore__ void foo(__gm__ int* buf) { *buf = 1; } int main(int argc, char* argv[]) { int a[100]; // nullptr - No L2 usage here foo<<<5, nullptr, nullptr>>>(a); return 0; }
编译命令如下:
# 功能:Host & Device代码混合编译,生成可执行文件 axpy # 编译命令: $bisheng --cce-soc-version=AscendXXXYY --cce-soc-core-type=VecCore -O2 axpy.cce -o axpy -I$RT_INC -L$RT_LIB
- 运行可执行文件。
# 运行: $export LD_LIBRARY_PATH=$RT_LIB:$LD_LIBRARY_PATH $./axpy