Compilation, Link, and Library Usage

Sample Code

This section uses the following executable sample as an example. Before using the sample, configure the environment first. For details, see Heterogeneous Compile Procedure.

// .aicpu file, which contains the AI CPU kernel function code.
// 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;
}

// Use the .cce file as an example. The AI CPU kernel is heterogeneously called using <<<>>>.
// 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);

  // Allocate the GM space, and pass the GM pointer as the input parameter to the AI CPU.
  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;
}

Generating and Using a Dynamic Link Library Through the AI CPU File

The compile and link commands are as follows. Replace ${INSTALL_DIR} with the file storage path after the CANN software is installed.

# Compile command:
# Compile the device file to generate the AI CPU binary file, with the AI CPU device code embedded. The options following test_aicpu.o in the compile command are those required for AscendC::printf.
$bisheng -O2 test.aicpu -c -o test_aicpu.o --cce-aicpu-L${INSTALL_DIR}/lib64/device/lib64 --cce-aicpu-laicpu_api --cce-aicpu-toolkit-path=${INSTALL_DIR}/toolkit/toolchain/hcc/bin -I${INSTALL_DIR}/include/ascendc/aicpu_api -D__AICPU_DEVICE__
# Compile and link the AI CPU binary file into a dynamic link library. libtest_aicpu.so is also a dynamic library on the host, with the device code embedded. (This step must be performed separately from the previous step.)
$bisheng --shared -o libtest_aicpu.so test_aicpu.o -L${INSTALL_DIR}/lib64 -lascendc_runtime
# Compile the host file into an executable file and link the dynamic link library during compile. libpthread.so, libstdc++.so, and libdl.so are the required standard libraries.
$bisheng test_main.cce -I$RT_INC -L$RT_LIB -lprofapi -lunified_dlog -lascendcl -lruntime -lc_sec -lmmpa -lerror_manager -lascend_dump -L ./ -ltest_aicpu -lpthread lstdc++ -ldl -o test

Generating and Using a Static Link Library Through the AI CPU File

The compile and link commands are as follows. Replace ${INSTALL_DIR} with the file storage path after the CANN software is installed.

# Compile command:
# Compile the device file to generate the AI CPU binary file, with the AI CPU device code embedded. The options following test_aicpu.o in the compile command are those required for AscendC::printf.
$bisheng -O2 test.aicpu -c -o test_aicpu.o --cce-aicpu-L${INSTALL_DIR}/lib64/device/lib64 --cce-aicpu-laicpu_api --cce-aicpu-toolkit-path=${INSTALL_DIR}/toolkit/toolchain/hcc/bin -I${INSTALL_DIR}/include/ascendc/aicpu_api -D__AICPU_DEVICE__
# Compile and link the AI CPU binary file into a static link library. libtest_aicpu.so is also a static library on the host, with the device code embedded. (This step must be performed separately from the previous step.)
$ar x ${INSTALL_DIR}/lib64/libascendc_runtime.a 
$ar rcs libtest_aicpu.a test_aicpu.o */.o
# Compile the host file into an executable file and link the static link library during compile. libpthread.so, libstdc++.so, and libdl.so are the required standard libraries.
$bisheng test_main.cce -I$RT_INC -L$RT_LIB -lprofapi -lunified_dlog -lascendcl -lruntime -lc_sec -lmmpa -lerror_manager -lascend_dump -L ./ -ltest_aicpu -lpthread lstdc++ -ldl -o test

Case Execution

The execution result in the preceding scenarios is as follows:

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