异构编译步骤

毕昇编译器支持AI CPU异构编译，规范如下：

• AI CPU Device函数定义：需定义在.aicpu文件中或通过-x aicpu指定语法。
• AI CPU Kernel异构调用：需要在.cce/.asc文件中完成或通过-x cce/asc指定语法。

也就是说，.cce/.asc文件不可包含AI CPU Device函数定义，仅允许AI Core Device函数的定义以及AI Core/AI CPU Kernel的<<<>>异构调用。

在.aicpu文件和.cce/.asc文件中编程完毕后，对两类文件分别编译后再进行链接。示意图如下：

当需要编写包含异构程序时，可按如下步骤进行：

1. 从CANN发布包中获取毕昇编译器程序（bisheng）的路径，并设置环境变量。${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。以root用户安装为例，则安装后文件存储路径为：/usr/local/Ascend/cann。

   # 设置编译器环境变量，如：
   $export PATH=${INSTALL_DIR}/compiler/ccec_compiler/bin:$PATH

2. 代码开发。编写AI CPU异构程序的基本流程如下：
   1. 创建.aicpu文件（或者在编译时通过-x aicpu指定语法）。
   2. 在aicpu文件中编写设备Kernel函数。
   3. 创建.cce/.asc文件（或者在编译时通过-x cce/asc指定语法）。
   4. 在cce文件中include必需的Runtime（运行时管理）头文件。
   5. 在cce文件中通过extern声明需要调用的AI CPU kernel函数。
   6. 通过Runtime API创建device和stream。
   7. 申请设备内存，按需将输入数据拷贝至设备内存。
   8. 通过<<<>>>内核调用符运行AI CPU设备Kernel函数。
   9. 按需将输出数据拷贝至主机内存，并释放设备内存。
3. 获取runtime的动态链接库以及头文件的位置。在本示例中，编译依赖和所在位置如下：

   编译依赖	内容	以下简称
   runtime动态链接库路径	${INSTALL_DIR}/lib64/	$RT_LIBPATH
   runtime动态链接库	具体需要链接的库请参考5.2.3 编译、链接、库使用方法。	$RT_LIB
   runtime头文件	${INSTALL_DIR}/include	$RT_INC

4. 编写并执行编译命令，通过分步编译生成可执行文件。下文场景示例为开发环境和运行环境一致的情况。本节仅展示整体编译流程，代码为无意义场景不可运行，可运行样例请参考编译、链接、库使用方法。
   foo.aicpu：包含AI CPU Device代码

   // 简化写法，无意义场景不可运行，仅演示编译使用方法
   __global__ __aicpu__ int foo_aicpu(void* buf)
   {
       return 0;
   }

   foo.cce：包含AI CPU <<<>>>调用

   // 简化写法，无意义场景不可运行，仅演示编译使用方法
   extern __global__ __aicpu__ int foo_aicpu(void* buf);
   
   //支持cce文件中存在aicore函数定义及调用
   __global__ __aicore__ void foo_aicore(__gm__ int* buf)
   {
       *buf = 1;
   }
   
   int main(int argc, char* argv[])
   {
     aclrtStream stream;
     aclrtCreateStream(&stream);
     int a[100];
     foo_aicpu<<<1, nullptr, stream>>>((void*)a， sizeof(a));
     foo_aicore<<<5, nullptr, stream>>>(a);
     return 0;
   }

   编译命令如下：

   # 功能：aicpu和cce文件分步编译，链接生成可执行文件foo
   # 编译命令：
   $bisheng -O2 foo.aicpu -c -o foo.aicpu.o 
   $bisheng --npu-arch=dav-2201 -O2 foo.cce -c -o foo.cce.o
   $bisheng foo.aicpu.o foo.cce.o -o foo -I$RT_INC -L$RT_LIBPATH -l$RT_LIB