Compilation Using BiSheng Command Lines

The BiSheng Compiler, designed for Ascend AI Processors, can support heterogeneous programming extension and compile Ascend operator code into binary executable files and dynamic libraries. The executable program of the BiSheng Compiler is named bisheng, which supports host systems such as x86 and AArch64. It also natively supports the compilation of AI Core architecture instruction sets on the device. By using the BiSheng Compiler, you can program and develop Ascend AI Processors more efficiently.

Example for Quick Start

The following is an example of using the BiSheng Compiler to compile the add_custom operator with a static shape. This example shows how to compile the source file add_custom.asc and the specific compilation command. With this example, you can understand how to use the BiSheng Compiler to compile operators. For the complete example, click LINK.

  1. Include header files.
    When compiling the operator source file, you need to include necessary header files.
    1
    2
    3
    // Header files
    #include "acl/acl.h"
    #include "kernel_operator.h"
    
  2. Implement the kernel function.
    • The kernel function supports templates.
    • The input parameters of the kernel function can be user-defined structures, such as the AddCustomTilingData structure defined by the user in the example.
     1
     2
     3
     4
     5
     6
     7
     8
     9
    10
    11
    12
    13
    14
    15
    16
    17
    18
    19
    20
    21
    // User-defined TilingData structure
    struct AddCustomTilingData { 
        uint32_t totalLength; 
        uint32_t tileNum;
    };
    
    // Kernel implementation logic, including movement and computation
    class KernelAdd {
    public:
        __aicore__ inline KernelAdd() {}
        // ...
    
    };
    
    __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling)
    {
        KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); // Only the Vector core on the AI Core is started when the operator is executed.
        KernelAdd op;
        op.Init(x, y, z, tiling.totalLength, tiling.tileNum); 
        op.Process();
    }
    
  3. Call the function logic on the host, including memory allocation and release, initialization and deinitialization, and kernel function call using the kernel launch symbol.
     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
    27
    28
    29
    30
    31
    32
    33
    34
    35
    36
    37
    38
    39
    40
    41
    42
    // Header file to be included in the application on the host
    #include "acl/acl.h"
    // Header file to be included in the kernel
    #include "kernel_operator.h"
    // Kernel function development
    ...
    
    __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling)
    {
        KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);    
        KernelAdd op;
        op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
        op.Process();
    }
    
    // Call the operator using the kernel launch symbol <<<...>>>.
    std::vector<float> kernel_add(std::vector<float> &x, std::vector<float> &y)
    {
    ...
    }
    
    
    // Compare the computation result.
    uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
    {
    ...
    }
    
    // Main program for operator verification
    int32_t main(int32_t argc, char *argv[])
    {
        constexpr uint32_t totalLength = 8 * 2048;
        constexpr float valueX = 1.2f;
        constexpr float valueY = 2.3f;
        std::vector<float> x(totalLength, valueX);
        std::vector<float> y(totalLength, valueY);
    
        std::vector<float> output = kernel_add(x, y);
    
        std::vector<float> golden(totalLength, valueX + valueY);
        return VerifyResult(output, golden);
    }
    
  4. Run the following command to perform compilation:
    • -o demo: specifies the output file name as demo.
    • --npu-arch=dav-2201: specifies the NPU architecture version as dav-2201. The number following dav- indicates the NPU architecture version. For details about the architecture version number of each product model, see Table 1.
    1
    bisheng add_custom.asc -o demo --npu-arch=dav-2201
    
  5. Run the executable file.
    1
    ./demo
    

Program Compilation and Execution

The BiSheng Compiler can compile operator source files (with the .asc extension) into executable files, dynamic operator libraries, or static operator libraries on the current platform. In addition, C++/C source files with the .cpp/.c extension can also be compiled, but the -x asc compilation option needs to be added.

  • Compiling and generating an executable file
    1
    2
    3
    # 1. Compile hello_world.cpp into an executable file on the current platform.
    # bisheng [Operator source file] -o [Output product name] --npu-arch=[NPU architecture version]. The common parameter sequence is the same as that of g++.
    bisheng -x asc add_custom.cpp -o add_custom --npu-arch=dav-xxxx
    

    The generated executable file can be executed using the following command:

    ./add_custom
  • Compile and generate a dynamic operator library.
    1
    2
    3
    4
    # 2. Compile add_custom_base.cpp to generate a dynamic operator library.
    # bisheng -shared [Operator source file] -o [Output product name] --npu-arch=[NPU architecture version]
    # Dynamic library
    bisheng -shared -x asc add_custom_base.cpp -o libadd.so --npu-arch=dav-xxxx
    
  • Compile and generate a static operator library.
    1
    2
    3
    4
    # 3. Compile add_custom_base.cpp to generate a static operator library.
    bisheng -lib [Operator source file] -o [Output product name] --npu-arch=[NPU architecture version]
    # Static library
    bisheng -lib -x asc add_custom_base.cpp -o libadd.a --npu-arch=dav-xxxx
    

In the command line compilation scenario, you can link the required library files as needed. For details about common library files, see Table 1. During compilation, the library files listed in Table 2 are linked by default. Note the following exception scenario: When using g++ to link the static library generated by compiling the ASC code, you need to manually link the default link library.