编译、链接、库使用方法

代码示例

本章节编译指导均以此组实际可运行样例为示例。使用前需要先配置环境，详见4.2.2 异构编译步骤。

// aicpu文件，包含AI CPU Kernel函数代码
// test.aicpu
#include"aicpu_api.h"
struct KernelArgs {
    unsigned int a;
    unsigned int b;
    unsigned int *Device;
};

unsigned int NormalFunc(unsigned int a, unsigned int b) {
  return a + b;
}

__global__ __aicpu__ unsigned int KernelAicpu(void* args) {
  KernelArgs* cfg = (KernelArgs*)args;
  *(cfg->Device) = NormalFunc(cfg->a, cfg->b);
  AscendC::printf("Device result %d\n", *(cfg->Device));
  return 0;
}

// 以cce文件为例，通过<<<>>>异构调用AI CPU Kernel
// test.cce
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

struct KernelArgs {
    unsigned int a;
    unsigned int b;
    unsigned int *Device;
};

#define SPACESIZE 4096

extern __global__ __aicpu__ unsigned int KernelAicpu(void* args);

int main(int argc, char *argv[]) {
  aclInit(nullptr);
  aclrtSetDevice(0);
  aclrtStream stream;
  aclrtCreateStream(&stream);

  // 开辟GM空间,并将GM指针当做入参传给aicpu
  void* OutputDevice = nullptr;
  aclrtMalloc((void **)&OutputDevice, SPACESIZE, ACL_MEM_MALLOC_HUGE_FIRST);

  void* InputHost= malloc(SPACESIZE);
  memset(InputHost, 0, SPACESIZE);
  aclrtMemcpy(OutputDevice, SPACESIZE, InputHost, 4096, ACL_MEMCPY_HOST_TO_DEVICE);

  KernelArgs args;
  args.Device = (unsigned int*) OutputDevice ;
  args.a = 1;
  args.b = 2;
  // Invoke a kernel
  KernelAicpu<<<1, nullptr, stream>>>(&args, sizeof(KernelArgs));
  aclrtSynchronizeStream(stream);

  uint8_t *OutputHost = nullptr;
  aclrtMallocHost((void **)&OutputHost, SPACESIZE );
  aclrtMemcpy(OutputHost, SPACESIZE, OutputDevice,
                   SPACESIZE, ACL_MEMCPY_DEVICE_TO_HOST);
  aclrtSynchronizeStream(stream);
  
  printf("Result: 0x%08x\n", *OutputHost);

  aclrtFreeHost(OutputHost);
  aclrtFree(OutputDevice);

  aclrtDestroyStream(stream);
  aclrtResetDevice(0);
  aclFinalize();
  return 0;
}

本示例一共2个文件，其中：
- test.aicpu中是纯AI CPU Device侧代码，并且调用了AscendC::printf接口演示格式化输出功能。
- test_main.cce是异构代码，对test.aicpu中定义的AI CPU Kernel通过<<<>>>调用。为了不影响篇幅，本示例未添加AI Core Kernel的异构调用。
为了不影响篇幅，本示例未使用单独头文件，使用头文件的场景跟其他类C语言相同。
此用例的开发环境和运行环境相同，即本地构建和运行。

AI CPU文件生成动态链接库并使用

编译链接命令如下，${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。

# 编译命令：
# 将Device文件编译生成AI CPU二进制，内嵌AI CPU Device代码，编译命令中test_aicpu.o之后为支持AscendC::printf所需选项
$bisheng -O2 test.aicpu -c -o test_aicpu.o --cce-aicpu-L${INSTALL_DIR}/toolkit/lib64/device/lib64 --cce-aicpu-laicpu_api --cce-aicpu-toolkit-path=${INSTALL_DIR}/toolkit/toolchain/hcc/bin -I${INSTALL_DIR}/include/ascendc/aicpu_api
# 将AI CPU二进制编译链接成动态链接库，libtest_aicpu.so亦为Host侧动态库，内嵌Device代码（与上一步必须分步进行）
$bisheng --shared -o libtest_aicpu.so test_aicpu.o -L${INSTALL_DIR}/lib64 -lascendc_runtime
# 将Host文件编译成可执行文件，编译中链接动态链接库，libpthread.so、libstdc++.so、libdl.so为所需标准库
$bisheng test_main.cce -I$RT_INC -L$RT_LIB -lprofapi -lascendalog -lascendcl -lruntime -lc_sec -lmmpa -lerror_manager -lascend_dump -L ./ -ltest_aicpu -lpthread lstdc++ -ldl -o test

AI CPU文件生成静态链接库并使用

编译链接命令如下，${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。

# 编译命令：
# 将Device文件编译生成AI CPU二进制，内嵌AI CPU Device代码，编译命令中test_aicpu.o之后为支持AscendC::printf所需选项
$bisheng -O2 test.aicpu -c -o test_aicpu.o --cce-aicpu-L${INSTALL_DIR}/toolkit/lib64/device/lib64 --cce-aicpu-laicpu_api --cce-aicpu-toolkit-path=${INSTALL_DIR}/toolkit/toolchain/hcc/bin -I${INSTALL_DIR}/include/ascendc/aicpu_api
# 将AI CPU二进制编译链接成静态链接库，libtest_aicpu.a亦为Host侧静态库，内嵌Device代码（与上一步必须分步进行）
$ar x ${INSTALL_DIR}/lib64/libascendc_runtime.a 
$ar rcs libtest_aicpu.a test_aicpu.o */.o
# 将Host文件编译成可执行文件，编译中链接静态链接库，libpthread.so、libstdc++.so、libdl.so为所需标准库
$bisheng test_main.cce -I$RT_INC -L$RT_LIB -lprofapi -lascendalog -lascendcl -lruntime -lc_sec -lmmpa -lerror_manager -lascend_dump -L ./ -ltest_aicpu -lpthread lstdc++ -ldl -o test

AI CPU的Device代码在编译当前文件时已经经过链接，因此多个AI CPU二进制文件下的Device内容不会再链接，需要保证当前AI CPU Device文件里的内容不依赖其他AI CPU Device文件。
为了使能profiling和调测等功能，需要在AI CPU算子编译链接时链接部分库：
- libascendc_runtime.a，Ascend C算子参数等组装库，需要与AI CPU Device文件编译生成二进制进行链接。
- libruntime.so，Runtime运行库。
- libprofapi.so，Ascend C算子运行性能数据采集库。
- libascendalog.so，CANN日志收集库。
- libmmpa.so，CANN系统接口库。
- libascend_dump.so，CANN维测信息库。
- libc_sec.so，CANN安全函数库。
- liberror_manager.so，CANN错误信息管理库。
- libascendcl.so，acl相关接口库。

用例运行

以上场景运行结果都为：

$./test
Get Result in host: 0x00000003
Device Result: 0x00000003

父主题： AI CPU编程指导