Customizing the Add Operator

This document guides you through the following tasks, so you can experience the basic process of Ascend C SIMT operator development.

  1. Analyze the operator to determine the mathematical expression and computation logic.
  2. Develop the kernel function of the Add operator.
  3. Run the kernel function for verification.

Before development, you need to set up the environment. The Ascend C operator development process is illustrated in the following figure.

Figure 1 Ascend C operator development process
  • Click LINK to obtain the sample code.
  • To use this guide, you only need to have basic knowledge of C or C++. You can deepen your theoretical understanding of the Ascend C SIMT programming model during practical operations based on the knowledge. If you have no idea about the Ascend C SIMT programming model, you can try to run the sample in the guide first, and further progress your study by referring to the instructions at the end of the guide.

Environment Setup

Operator Analysis

Analyze the mathematical expression, number of inputs and outputs, shape range, and computation logic implementation of an operator, and specify the Ascend C SIMT APIs or operators to be called. The following uses the Add operator as an example to describe the analysis process.

  1. Specify the mathematical expression and computation logic of an operator.

    The Add operator is used as an example.

    The computation logic is as follows: Add inputs x and y at the corresponding positions in the global memory element-wise, and store the result in the global memory output z.

  2. Specify the input and output.

    • The Add operator has two inputs, x and y, and outputs the result z.
    • The supported input data type is float, and so is the output data type.
    • The input and output shapes of the operator are (48, 256).

  3. Define the kernel function name and parameters.

    • In this sample, the kernel function is named add_custom.
    • Based on the operator input and output analysis, it is determined that the kernel function has three input and output parameters: x, y, and z. All these parameters are of the float type.
    • Add an input parameter total_length to the kernel function to record the actual input and output data lengths of the operator. The data type is uint64_t.

  4. Determine the operator implementation logic.

    • Evenly distribute the data across 48 thread blocks. Each thread block starts 256 threads to process 256 elements, with each thread processing one element.
    • Calculate the offset of the data to be processed by the current thread based on the unique index of each thread.

Based on the preceding analysis, the design specifications of the Add operator implemented using Ascend C SIMT are as follows.
Table 1 Input and output specifications of the Add operator

Name

Shape

Data Type

Format

x (input)

48 × 256

float*

ND

y (input)

48 × 256

float*

ND

z (output)

48 × 256

float*

ND

total_length

-

uint64_t

-

Kernel Function Development

Compute the index of the current thread based on blockIdx (current thread block index), blockDim (number of threads in a single thread block), and threadIdx (current thread index). Then, use the computed index of the current thread as the offset of the data line currently being computed.
1
int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
Use the subscript offset and addition operator to compute the sum of the data at the offset position, and write the result to the output.
1
z[idx] = x[idx] + y[idx];

The complete kernel function code is as follows:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
__global__ void add_custom(float* x, float* y, float* z, uint64_t total_length)
{
    // Calculate global thread ID
    int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    // Maps to the row index of output tensor
    if (idx >= total_length) {
        return;
    }
    z[idx] = x[idx] + y[idx];
}

Runtime Verification of the Kernel Function

After the kernel function is developed on the kernel, you can compile the kernel function calling program on the host, implementing the function of calling operators from the application program on the host to execute the computation process.

  1. Compile the program framework on the host.

     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
47
48
    // Header file required for host calling
#include <vector>
#include "acl/acl.h"

// Kernel function development
__global__ void add_custom(float* x, float* y, float* z, uint64_t total_length)
{
    ...
}

// Call the operator using the kernel launch symbol <<<...>>>.
std::vector<float> add(std::vector<float>& x, std::vector<float>& y)
{
    ...
    // Calc splite params
    uint32_t block_num = 48;
    uint32_t thread_num_per_block = 256;
    uint32_t dyn_ubuf_size = 0;  // No need to alloc dynamic memory.
    // Call kernel function with <<<...>>>
    add_custom<<<block_num, thread_num_per_block, dyn_ubuf_size, stream>>>(x_device, y_device, z_device, x.size());
    ...
    return output;
}

// Compare the computation result.
uint32_t verify_result(std::vector<float>& output, std::vector<float>& golden)
{
    if (std::equal(output.begin(), output.end(), golden.begin())) {
        std::cout << "[Success] Case accuracy is verification passed." << std::endl;
        return 0;
    } else {
        std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
        return 1;
    }
    return 0;
}

// Verify the operator main program.
int32_t main(int32_t argc, char* argv[])
{
    constexpr uint32_t in_shape = 48 * 256;
    std::vector<float> x(in_shape);
    std::vector<float> y(in_shape);
    std::vector<float> golden(in_shape);
    ...
    std::vector<float> output = add(x, y);
    return verify_result(output, golden);
}

  2. Write the code for calling the operator using the kernel launch symbol <<<...>>>.

    Figure 2 Calling process

    For details about how to use the acl APIs in the following example, see Runtime API Reference.

     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
    std::vector<float> add(std::vector<float>& x, std::vector<float>& y)
{
    size_t total_byte_size =x.size() * sizeof(float);
    int32_t device_id = 0;
    aclrtStream stream = nullptr;
    uint8_t* x_host = reinterpret_cast<uint8_t *>(x.data());
    uint8_t* y_host = reinterpret_cast<uint8_t *>(y.data());
    uint8_t* z_host = nullptr;
    float* x_device = nullptr;
    float* y_device = nullptr;
    float* z_device = nullptr;
    // Init
    aclInit(nullptr);
    aclrtSetDevice(device_id);
    aclrtCreateStream(&stream);
    // Malloc memory in host and device
    aclrtMallocHost((void **)(&z_host), total_byte_size);
    aclrtMalloc((void **)&x_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&y_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMalloc((void **)&z_device, total_byte_size, ACL_MEM_MALLOC_HUGE_FIRST);
    aclrtMemcpy(x_device, total_byte_size, x_host, total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
    aclrtMemcpy(y_device, total_byte_size, y_host, total_byte_size, ACL_MEMCPY_HOST_TO_DEVICE);
    // Calc splite params
    uint32_t block_num = 48;
    uint32_t thread_num_per_block = 256;
    uint32_t dyn_ubuf_size = 0;  // No need to alloc dynamic memory.
    // Call kernel function with <<<...>>>
    add_custom<<<block_num, thread_num_per_block, dyn_ubuf_size, stream>>>(x_device, y_device, z_device, x.size());
    aclrtSynchronizeStream(stream);
    // Copy result from device to host
    aclrtMemcpy(z_host, total_byte_size, z_device, total_byte_size, ACL_MEMCPY_DEVICE_TO_HOST);
    std::vector<float> output((float *)z_host, (float *)(z_host + total_byte_size));
    // Free memory
    aclrtFree(x_device);
    aclrtFree(y_device);
    aclrtFree(z_device);
    aclrtFreeHost(z_host);
    // DeInt
    aclrtDestroyStream(stream);
    aclrtResetDevice(device_id);
    aclFinalize();
    return output;
}

  3. Configure CMake compilation as follows:

     1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
    cmake_minimum_required(VERSION 3.16)
# find_package(ASC) is a command used in CMake to search for and configure the Ascend C compilation toolchain.
find_package(ASC REQUIRED)
# Specify that the project supports the ASC and CXX languages. ASC indicates that the Ascend C programming language can be compiled using the BiSheng compiler.
project(kernel_samples LANGUAGES ASC CXX)

add_executable(demo
    add.asc
)

# Set the NPU architecture through the compilation option.
target_compile_options(demo PRIVATE   
   $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510 --enable-simt>
)

  4. Execute the compilation and running commands.

    1
2
3
    mkdir -p build && cd build; 
cmake ..;make -j;
./demo
    • This sample supports only the following models:
      • Atlas 350 Accelerator Card
    • --enable-simt specifies the SIMT programming scenario.
    • --npu-arch is used to specify the architecture version of the NPU. The string following dav- is the architecture version number. You need to replace the architecture version number with the actual one. For details about the mapping between the AI processor model and architecture version number, see Table 1.

Follow-Up Guide

To learn more about SIMT programming concepts, you can refer to AI Core SIMT Programming to learn the basic concepts and then review this tutorial. If you already understand related concepts and can run the sample through, refer to SIMT Operator Implementation for more details about Ascend C SIMT programming.

Parent topic: SIMT-based Programming