编译、链接、库使用方法

代码示例

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

// 一个Device侧文件，包含Kernel函数代码
// test.cce
__global__ [aicore] void foo(__gm__ uint8_t *Out, int Stride) {
  Out[block_idx * Stride] = block_idx;
}

// 一个Host侧文件，通过<<<>>>异构调用Kernel代码 
// test_main.cce
#include "acl/acl.h"
#include <stdio.h>
#include <stdlib.h>

#define BLOCKS 4
#define CACHELINE_SZ 64

extern void __global__ [aicore] foo(__gm__ uint8_t *Out, int Stride);

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

  uint8_t ExpectedValue[] = {0, 1, 2, 3};
  uint8_t *OutputValue = nullptr;
  aclrtMalloc((void **)&OutputValue, BLOCKS, ACL_MEM_MALLOC_NORMAL_ONLY);

  uint8_t InitValue[BLOCKS] = {0};
  aclrtMemcpyAsync((void *)OutputValue, sizeof(InitValue), InitValue,
                   sizeof(InitValue), ACL_MEMCPY_HOST_TO_DEVICE, stream);
  aclrtSynchronizeStream(stream);

  // Invoke a kernel
  foo<<<BLOCKS, nullptr, stream>>>(OutputValue, CACHELINE_SZ);

  uint8_t *OutHost = nullptr;
  aclrtMallocHost((void **)&OutHost, BLOCKS * CACHELINE_SZ);
  aclrtMemcpyAsync(OutHost, BLOCKS * CACHELINE_SZ, OutputValue,
                   BLOCKS * CACHELINE_SZ, ACL_MEMCPY_DEVICE_TO_HOST, stream);
  aclrtSynchronizeStream(stream);

  for (int I = 0; I < sizeof(ExpectedValue) / sizeof(uint8_t); I++) {
    printf("i%d\t Expect: 0x%04x\t\t\t\tResult: 0x%04x\n", I, ExpectedValue[I],
           OutHost[I * CACHELINE_SZ]);
  }

  aclrtFreeHost(OutHost);
  aclrtFree(OutputValue);

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

本示例一共2个文件，其中，test.cce是纯Device侧代码，test_main.cce是Host异构代码。这是一组比较典型，且能覆盖大部分场景的例子。为了不影响篇幅，本示例未使用单独头文件，使用头文件的场景跟其他类C语言相同。

此用例的开发环境和运行环境相同，即本地构建和运行。

生成动态链接库并使用

# 编译命令：
# 将device文件编译链接成动态链接库，libabcd.so亦为host侧动态库，内嵌device代码
$bisheng -O2 -fPIC --npu-arch=dav-2201 test.cce --shared -o libabcd.so
# 将host文件编译成可执行文件，编译中链接动态链接库
$bisheng test_main.cce -I$RT_INC -L$RT_LIB -lascendcl -L ./ -labcd -lruntime -o test

生成静态链接库并使用

将Device文件编译成静态链接库后，使用bisheng链接该静态库并把Host文件编译成可执行文件。

# 编译命令：
# 将device文件编译链接成libstaticabcd.a静态链接库
$bisheng -O2 -fPIC --npu-arch=dav-2201 -I$RT_INC test.cce --cce-build-static-lib  -o libstaticabcd.a
# 将host文件编译成可执行文件，编译中链接静态链接库
$bisheng test_main.cce -I$RT_INC -L$RT_LIB -L ./ -lstaticabcd  -lruntime -lascendcl -lstdc++ -o test

使用--cce-build-static-lib可以直接将多文件编译并打包成Host侧静态库，内嵌已链接的Device代码。请特别注意，因为这个编译模式下的Device代码已经经过链接，不适合再与其他内嵌可重定位device obj的目标文件链接。

用例运行

以上场景运行结果都为：

$./test
i0	 Expect: 0x0000				Result: 0x0000
i1	 Expect: 0x0001				Result: 0x0001
i2	 Expect: 0x0002				Result: 0x0002
i3	 Expect: 0x0003				Result: 0x0003