目录

Kernel Function

According to the SPMD model, when using Ascend C for programming, we write an operator implementation code. When the operator is called, N running instances are started and run on N cores. This section describes the entry function for implementing an operator.

The kernel function is an entry for implementing an Ascend C operator on the device. In a kernel function, to-be-performed data access and operations need to be specified for code executed on a core. When the kernel function is called, multiple cores will execute the same kernel function code in parallel, with the same parameters.

Ascend C allows you to use kernel functions, which are syntax extensions of C/C++ functions, to manage the running code on the device. You can create operator objects and call their member functions in kernel functions to implement all functions of the operator. The kernel function is a bridge between the host and the device. This chapter describes how to use the kernel function.

Kirin 9020, Kirin 9030, and Kirin X90 series processors do not support kernel function calls using <<<>>>.

Kernel Function Definition

extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z)

The preceding is an example of kernel function declaration. Comply with the following rules when writing kernel functions:

Kernel Function Call

The call statement of the kernel function is an extension of the C/C++ function call statement. This section describes only the basic calling method. During actual operator development, you can call operators using this method or by connecting to the Ascend C framework. For details, see Project-based Operator Development.

The common function calling mode is as follows:

function_name(argument list);

The kernel function uses the kernel launch symbol <<<...>>> to specify the execution configuration of the kernel function.

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

The kernel launch symbol can be called only when performing compilation on the NPU. It cannot be identified by compilation on the CPU.

The execution configuration is determined by the following blockDim, l2ctrl, and stream parameters:

[h2]blockDim

Specifies the number of cores on which a kernel function will be executed. A logical ID, that is, block_idx, is allocated to each core that executes the kernel function, and block_idx can be obtained by calling GetBlockIdx during the kernel function implementation.

blockDim is a concept about logical cores, and its value range is [1, 65535]. To fully utilize hardware resources, set this parameter to the number of physical cores or a multiple of the number of physical cores. For the coupled architecture and decoupled architecture, the meaning and setting rules of blockDim during running are different. The details are as follows:

[h2]l2ctrl

Reserved parameter and can be ignored. This parameter is fixed at nullptr.

[h2]stream

aclrtStream type. Streams preserve the order of some asynchronous operations being executed on the device.

[h2]Calling Example

The following kernel function add_custom is used to add two vectors. The following is a calling sample:

// If blockDim is set to 8, the add_custom kernel function is called on eight cores. Each core independently executes the kernel function in parallel. The argument lists of the kernel function are x, y, and z.
add_custom<<<8, nullptr, stream>>>(x, y, z);

The kernel function is called asynchronously. After the kernel function is called, the control right is returned to the host immediately. You can call the aclrtSynchronizeStream function to force the host program to wait until all kernel functions are executed.

aclError aclrtSynchronizeStream(aclrtStream stream);

Kernel Function Sample

The following provides sample code snippets for implementing kernel functions. For details about the complete sample, see Vector Programming.

// Implement the kernel function.
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
    // Initialize the operator class. The operator class provides methods such as operator initialization and core processing.
    KernelAdd op;
    // Initialization function, which is used to obtain the input and output addresses to be processed by the kernel function and initialize the memory.
    op.Init(x, y, z);
    // Core processing function, which is used to complete core logic such as data movement and computation of operators.
    op.Process();
}