AI Core Custom Operator Loading

This section describes how to load and execute a custom kernel developed by using Ascend C. You need to call the APIs described in this section only when you write operator kernels or need to directly control the kernel binary loading, argument packing, and task dispatch. If no operators are written and only the operators provided by CANN are called, the aclnn operator APIs, such as aclnnAdd, should be called first. The aclnn APIs encapsulate the operator selection, argument processing, and kernel loading and execution processes. Typically, you do not need to directly call the LaunchKernel API.

There are mainly two modes for launching a custom kernel:

• Hybrid programming: uses the <<< >>> syntax in the host code to launch the kernel.
• Non–hybrid programming: compiles the kernel into a standalone operator binary, and then loads and launches the binary using Runtime APIs such as aclrtBinaryLoadFromFile, aclrtBinaryGetFunction, and aclrtLaunchKernelWithConfig.

1
2
3
4
5

include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)
ascendc_library(kernels STATIC kernel_print.cpp)

add_executable(main main.cpp)
target_link_libraries(main PRIVATE kernels ${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so)

The host code generated in the preceding mode can directly call the kernel launch function. You do not need to explicitly call APIs such as aclrtBinaryLoadFromFile and aclrtBinaryGetFunction.

The following is a key code example. It is for reference only and cannot be directly copied for compilation or running:

 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
43
44
45
46

// Device code
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

int main()
{
    int64_t n = ...;
    size_t size = static_cast<size_t>(n) * sizeof(uint64_t);

    aclInit(nullptr);
    aclrtSetDevice(0);

    aclrtStream stream = nullptr;
    aclrtCreateStream(&stream);

    void *hX = nullptr;
    void *hY = nullptr;
    void *hZ = nullptr;
    aclrtMallocHost(&hX, size);
    aclrtMallocHost(&hY, size);
    aclrtMallocHost(&hZ, size);

    // Initialize the input data.
    ...

    void *dX = nullptr;
    void *dY = nullptr;
    void *dZ = nullptr;
    aclrtMalloc(&dX, size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&dY, size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&dZ, size, ACL_MEM_MALLOC_HUGE_FIRST);

    aclrtMemcpy(dX, size, hX, size, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(dY, size, hY, size, ACL_MEMCPY_HOST_TO_DEVICE);

    uint32_t numBlocks = 48;
    add_custom<<<numBlocks, nullptr, stream>>>(dX, dY, dZ);
    aclrtSynchronizeStream(stream);

    aclrtMemcpy(hZ, size, dZ, size, ACL_MEMCPY_DEVICE_TO_HOST);
    ...
}

1
2
3
4
5

include(${ASCENDC_CMAKE_DIR}/ascendc.cmake)
ascendc_fatbin_library(ascendc_kernels_simple add_custom.cpp)

add_executable(ascendc_kernels_bbit main.cpp)
target_link_libraries(ascendc_kernels_bbit PRIVATE ${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so)

The generated operator binary file is loaded at runtime via a path, for example, ./out/fatbin/ascendc_kernels_simple/ascendc_kernels_simple.o. In this mode, the host code does not directly reference the kernel function symbols. Instead, it operates the kernel through the binary and function handles.

The related concepts are as follows:

• Binary: a dynamically loaded code container unit, which contains the compiled kernel code and global variables. You can call aclrtBinaryLoadFromFile or aclrtBinaryLoadFromData to load the operator binary and get the binary handle.
• Function: a specific, executable kernel entry point within a binary. You can call aclrtBinaryGetFunction or aclrtBinaryGetFunctionByEntry to get the function handle.
• Argument list: The LaunchKernel API needs to get kernel arguments. The arguments can be stored in device memory, host memory, or the aclrtArgsHandle argument list. You can also use placeholders to enable Runtime to move small blocks of argument data during launch.

The following is a key code example for using the LaunchKernel API. It is for reference only and cannot be directly copied for compilation or running:

// Device code, which is compiled into a standalone operator binary file.
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    KernelAdd op;
    op.Init(x, y, z);
    op.Process();
}

int main()
{
    int64_t n = ...;
    size_t size = static_cast<size_t>(n) * sizeof(uint64_t);

    aclInit(nullptr);
    aclrtSetDevice(0);

    aclrtStream stream = nullptr;
    aclrtCreateStream(&stream);

    // Load the operator binary and get the kernel function handle.
    aclrtBinHandle bin = nullptr;
    aclrtBinaryLoadFromFile("add_custom.o", nullptr, &bin);

    aclrtFuncHandle func = nullptr;
    aclrtBinaryGetFunction(bin, "add_custom", &func);

    void *hX = nullptr;
    void *hY = nullptr;
    void *hZ = nullptr;
    aclrtMallocHost(&hX, size);
    aclrtMallocHost(&hY, size);
    aclrtMallocHost(&hZ, size);

    // Initialize the input data.
    ...

    void *dX = nullptr;
    void *dY = nullptr;
    void *dZ = nullptr;
    aclrtMalloc(&dX, size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&dY, size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc(&dZ, size, ACL_MEM_MALLOC_HUGE_FIRST);

    aclrtMemcpy(dX, size, hX, size, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(dY, size, hY, size, ACL_MEMCPY_HOST_TO_DEVICE);

    uint32_t numBlocks = 48;
    void *args[] = {dX, dY, dZ};
    size_t argsSize = sizeof(args);
    aclrtLaunchKernelWithHostArgs(func, numBlocks, stream, nullptr, args, argsSize, nullptr, 0);
    aclrtSynchronizeStream(stream);

    aclrtMemcpy(hZ, size, dZ, size, ACL_MEMCPY_DEVICE_TO_HOST);

    aclrtBinaryUnLoad(bin);
    ...
}

The following are recommendations on selecting a LaunchKernel API.

For a complete sample, see 0_launch_kernel. This sample demonstrates the process of binary loading, function retrieval, aclrtArgsHandle argument packing, placeholder arguments, and kernel launch via aclrtLaunchKernelWithConfig.