AI CPU Operator Compilation

Compilation Using BiSheng Command Lines

The following uses a "Hello World" print example to describe how to compile an AI CPU operator using BiSheng command lines.

The content of the hello_world.aicpu file is as follows:

1
2
3
4
5
6
7

#include "aicpu_api.h"

__global__ __aicpu__ uint32_t hello_world(void *args)
{
    AscendC::printf("Hello World!!!\n");
    return 0;
}

On the host, the kernel launch symbol <<<...>>> is used to call the AI CPU operator. The sample code of main.asc is as follows:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

#include "acl/acl.h"

struct KernelArgs {
    int mode;
};

extern __global__ __aicpu__ uint32_t hello_world(void *args);

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

    struct KernelArgs args = {0};
    constexpr uint32_t blockDim = 1;
    hello_world<<<blockDim, nullptr, stream>>>(&args, sizeof(KernelArgs));
    aclrtSynchronizeStream(stream);

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

You can use BiSheng command lines to compile hello_world.aicpu and main.asc into .o files and then link them into executable files. The compilation commands are as follows:

• When compiling hello_world.aicpu, use -I to specify the path of the dependent header file, use --cce-aicpu-laicpu_api to link the dependent library libaicpu_api.a for the device, and use --cce-aicpu-L to specify the path of libaicpu_api.a.
• When compiling main.asc, use the --npu-arch option to specify the corresponding architecture version number.

Replace ${INSTALL_DIR} with the CANN component directory. For example, if the installation is performed by the root user, the default file storage path is /usr/local/Ascend/cann.

$bisheng -O2 hello_world.aicpu --cce-aicpu-L${INSTALL_DIR}/lib64/device/lib64 --cce-aicpu-laicpu_api -I${INSTALL_DIR}/include/ascendc/aicpu_api -c -o hello_world.aicpu.o
$bisheng --npu-arch=dav-2201 main.asc -c -o main.asc.o
$bisheng hello_world.aicpu.o main.asc.o -o demo

The preceding example demonstrates how to use BiSheng command lines to compile and generate an executable file. In addition, BiSheng command lines can be used to compile and generate dynamic and static libraries for AI CPU operators. You can use the kernel call symbol <<<...>>> in the ASC code to call the kernel function of an AI CPU operator. When compiling the ASC source file to generate an executable file, link the AI CPU dynamic or static library. Note: If you compile the AI CPU operator code separately to generate a dynamic or static library, you need to manually link default link libraries listed in Table 2.

• Compile and generate a dynamic operator library.

  # Compile test_aicpu.cpp to generate a dynamic operator library.
  # -lxxx indicates the default link library.
  # bisheng -shared -x aicpu test_aicpu.cpp -o libtest_aicpu.so -lxxx ...

• Compile and generate a static operator library.

  # Compile test_aicpu.cpp to generate a static operator library.
  # -lxxx indicates the default link library.
  # bisheng -lib -x aicpu test_aicpu.cpp -o libtest_aicpu.a -lxxx ...

Common Compilation Options for AI CPU Operators

Common compilation options of AI CPU operators are as follows.

Compilation Using CMake

In a project, you can use CMake to more easily use the BiSheng Compiler to compile AI CPU operators and generate executable files, dynamic libraries, static libraries, or binary files.

The "Hello World" print sample described in Compilation Using BiSheng Command Lines is still used as an example. In addition to the code implementation file, you need to prepare a CMakeLists.txt file in the project directory.

├── hello_world.aicpu // AI CPU operator kernel function definition
├── main.asc // AI CPU operator kernel function call
└── CMakeLists.txt

The content of CMakeLists.txt is as follows:

cmake_minimum_required(VERSION 3.16)
# 1. find_package() is a command used in CMake to search for and configure the Ascend C compilation toolchain.
find_package(ASC REQUIRED) 
find_package(AICPU REQUIRED) 

# 2. Specify that the project supports the ASC, AI CPU, and CXX languages. ASC indicates that the BiSheng Compiler can be used to compile the Ascend C programming language. AI CPU indicates that the BiSheng Compiler can be used to compile AI CPU operators.
project(kernel_samples LANGUAGES ASC AICPU CXX)

# 3. Use the CMake API to compile executable files.
add_executable(demo
    hello_world.aicpu
    main.asc
)

# 4. Specify the linker because the ASC and AI CPU languages are used.
set_target_properties(demo PROPERTIES LINKER_LANGUAGE ASC)  // Specify the language used for the link.

target_include_directories(demo PUBLIC
    ${INSTALL_DIR}/include/ascendc/aicpu_api    
)

target_compile_options(demo PRIVATE
    # --npu-arch specifies the NPU architecture version. dav- is followed by the architecture version number. For details about the architecture version number of each product model, see Table 1.
    # <COMPILE_LANGUAGE:ASC>: indicates that the compilation option takes effect only for the ASC language.
    $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
)

If you need to use CMake to compile and generate dynamic and static libraries, the following provides more detailed compilation examples:

• Compile the .cpp file to generate a dynamic library.

  # Set the .cpp file to the ASC attribute and enable the Ascend C language for compilation.
  set_source_files_properties(
      add_custom_base.cpp 
      sub_custom_base.cpp
      PROPERTIES LANGUAGE ASC
  )
  
  # Set the .cpp file to the AI CPU attribute to support AI CPU operator compilation.
  set_source_files_properties(
      aicpu_kernel.cpp
      PROPERTIES LANGUAGE AICPU
  )
  
  add_library(kernel_lib SHARED
      add_custom_base.cpp 
      sub_custom_base.cpp
      aicpu_kernel.cpp # AI CPU operators and AI Core operators can be packed together into a dynamic library.
  )
  # During AI CPU operator compilation, you need to manually link the following dependency libraries. (If the specified link language is ASC, you do not need to manually link the following libraries.)
  target_link_libraries(kernel_lib PRIVATE
      ascendc_runtime
      profapi
      ascendalog
      ascendcl
      runtime
      c_sec
      mmpa
      error_manager
      ascend_dump
  )
  
  add_executable(demo
      main.cpp
  )
  target_compile_definitions(demo PRIVATE
      ASCENDC_DUMP=0
  )
  target_compile_options(demo PRIVATE
      -g
  )
  target_include_directories(demo PRIVATE
      include
  )
  target_link_libraries(demo PRIVATE
      kernel_lib
  )

• Compile the .asc and .aicpu files to generate a static library.

  # By default, the Ascend C language is used to compile the .asc file, and the AI CPU language is used to compile the .aicpu file. You do not need to use set_source_files_properties to set the language.
  add_library(kernel_lib STATIC
      add_custom_base.asc 
      sub_custom_base.asc
      aicpu_kernel.aicpu # AI CPU operators and AI Core operators can be packed together into a static library.
  )
  
  add_executable(demo
      main.cpp
  )
  target_compile_definitions(demo PRIVATE
      ASCENDC_DUMP=0
  )
  target_compile_options(demo PRIVATE
      -g
  )
  target_include_directories(demo PRIVATE
      include
  )
  target_link_libraries(demo PRIVATE
      kernel_lib
  )

The following lists the common variable configurations and link libraries used during CMake compilation.

set(CMAKE_CXX_COMPILER_LAUNCHER <launcher_program>)

set(CMAKE_CXX_COMPILER_LAUNCHER /usr/bin/ccache)