昇腾社区首页
中文
注册

异构编译

编译选项

常用的编译选项说明如下:

选项

是否必需

说明

-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。

  • VecCore:编写的代码,仅包含Vector侧代码;
  • CubeCore:编写的代码,仅包含Cube侧代码;
  • AICore:编写的代码,同时包含Vector和Cube侧代码。

异构编译

毕昇编译器支持异构编译,当需要编写包含异构程序时,可按如下步骤进行:

  1. 从CANN发布包中获取毕昇编译器程序(bisheng)的路径,并设置环境变量。
    # 设置编译器环境变量,如:
    $export PATH=${INSTALL_DIR}/compiler/ccec_compiler/bin:$PATH
  2. 代码开发。编写异构程序的基本流程如下:
    1. include必需的AscendCL Runtime(运行时管理)头文件。
    2. 通过Runtime API创建devicestream
    3. 申请设备内存,按需将输入数据拷贝至设备内存。
    4. 编写设备kernel函数。
    5. 通过<<<>>>异构调用接口运行设备kernel函数。
    6. 按需将输出数据拷贝至主机内存,并释放设备内存。
  3. 获取runtime的动态链接库以及头文件的位置。在本示例中,编译依赖和所在位置如下:

    编译依赖

    位置

    以下简称

    runtime动态链接库

    ${INSTALL_DIR}/lib64/

    $RT_LIB

    runtime头文件

    ${INSTALL_DIR}/include

    $RT_INC

  4. 编写并执行编译命令,一步编译生成可执行文件。下文场景示例为开发环境和运行环境一致的情况。
    // 简化写法,无意义场景不可运行,仅演示编译使用方法
    #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
  5. 运行可执行文件。
     # 运行: 
     $export LD_LIBRARY_PATH=$RT_LIB:$LD_LIBRARY_PATH
     $./axpy