Kernel Function

Defining and Calling Kernel Functions

Comply with the following rules when defining kernel functions:

• Use function type qualifiers.

  In addition to defining kernel functions in C/C++ function declaration mode, you need to add additional function type qualifiers, including __global__ and __aicore__, to kernel functions.

  The __global__ function type qualifier is used to identify a kernel function that can be called by <<<...>>>. The __aicore__ function type qualifier is used to specify that the kernel function is executed on the AI Core of the device.

  1	__global__ __aicore__ void kernel_name(argument list);

  Functions used in programming can be classified into three types: kernel functions (executed on the device), host functions, and device functions (except kernel functions). The following figure uses operator development based on the kernel launch project as an example to describe the calling relationships between the three types of functions.

  • Host functions can call host functions of the same type, similar to function calling in general C/C++ programming. Kernel functions can also be called by <<<...>>>.
  • Device functions (except kernel functions) can call device functions of the same type.
  • Kernel functions can call device functions (except kernel functions).
  Figure 1 Calling relationships between kernel functions, host functions, and device functions
  
• Use variable type qualifiers.

  The variable type qualifier __gm__ needs to be added to the pointer input parameter variable, indicating that the pointer variable points to a memory address in the Global Memory.

• Other rules or suggestions
  1. Rule: The kernel function must have the void return type.
  2. Rule: The input parameter must be a pointer or C/C++ built-in data type (primitive data types), for example, half* s0, float* s1, and int32_t c.
  3. Suggestion: To unify the expression, you are advised to use the GM_ADDR macro to modify the input parameter. The definition of the GM_ADDR macro is as follows:

     1	#define GM_ADDR __gm__ uint8_t*

     An example of using GM_ADDR to modify the input parameter is as follows:

     1	extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)

     The pointer of the uint8_t type is used as an example. In subsequent use, the pointer needs to be converted into the actual pointer type.

The common function calling mode is as follows:

1

function_name(argument list);

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

1

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 following kernel function add_custom is used to add two vectors. The following is a calling sample:

1
2

//If blockDim is set to 8, the add_custom kernel function is called on 8 cores. Each core independently executes the kernel function in parallel. The parameter 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 aclrtSynchronizeStream to force the host program to wait until all kernel functions are executed.

1

aclError aclrtSynchronizeStream(aclrtStream stream);

Defining and Calling the Template Kernel Function

You can use a template to define a kernel function. The following is an example of defining a kernel function, which has two template parameters: a and T. a is a non-type template parameter, while T is a type template parameter.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11

template<int a, typename T>
__global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{
...
    AscendC::printf("Print Template a: %d\n", a);
...
    xGm.SetGlobalBuffer((__gm__T*)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
    yGm.SetGlobalBuffer((__gm__T*)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
    zGm.SetGlobalBuffer((__gm__T*)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
...
}

The template kernel function is called as follows: add_custom<20, float>. This part of code calls the kernel function named add_custom and provides specific values for its template parameters.

1

add_custom<20, float><<<blockDim, nullptr, stream>>>(x, y, z);