Documentation
Download

C API Programming Based on the Language Extension Layer

C API programming based on the language extension layer provides pure C-style APIs, complying with the development habits of C language operators and offering a programming experience similar to that in the industry. This section describes the C API programming paradigm, including memory management, synchronization control, and compute and movement APIs, to help developers better understand and use C APIs for programming.

Memory Management

C APIs use C-style address qualifiers to describe different levels of memory and can directly operate memory addresses through pointers, thereby precisely controlling the data storage location. The following table describes the address qualifiers for different storage units.

Table 1 Address qualifiers for different storage units

Storage Unit

Address Qualifier

Description

Global Memory

__gm__

Indicates that the modified variable is located in the global memory address space.

Unified Buffer

__ubuf__

Indicates that the modified variable is located in the unified buffer address space.

L1 Buffer

__cbuf__

Indicates that the modified variable is located in the L1 buffer address space.

L0A Buffer

__ca__

Indicates that the modified variable is located in the L0A buffer address space.

L0B Buffer

__cb__

Indicates that the modified variable is located in the L0B buffer address space.

L0C Buffer

__cc__

Indicates that the modified variable is located in the L0C buffer address space.

Address space qualifiers can be used in array or pointer variable declarations to specify the region in which an object is allocated. Multiple address space qualifiers cannot be used for the same type.

During programming based on C APIs, you need to control the memory through explicit memory management. The following describes how to allocate memory at different levels:

  • Global memory: Generally, the memory is allocated by calling the on-device aclrtMalloc API. The corresponding address qualifier needs to be added.
  • Internal storage (including the unified buffer and L1 buffer): The space is allocated by the user and declared in the kernel using the address qualifier keyword. There is no automatic garbage collection mechanism. You need to strictly control the lifecycle. The following uses UB space allocation as an example:
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
// Use the address space qualifier in the array variable declaration.
// total_length indicates the length of the data involved in computation.
constexpr uint64_t total_length = 256;
__ubuf__ float xLocal[ total_length ];
__ubuf__ float yLocal[ total_length ];
__ubuf__ float zLocal[ total_length ];

// Use the address space qualifier in the pointer variable declaration.
uint64_t offset = 0;                                   // Allocate memory for src0, starting from 0.
__ubuf__ half* src0 = (__ubuf__ half*)asc_get_phy_buf_addr(offset);    // Obtain the address of src0 and use the __ubuf__ keyword to specify that the address points to the UB memory.

Synchronization Control

The NPU has different compute units. Before computation, the data usually needs to be moved to the compute units. The computation and data movement processes on different compute units can be divided into different pipelines, as described in the following table.

Table 2 Instruction pipeline types

Pipeline Type

Description

PIPE_S

Scalar pipeline.

PIPE_V

Vector computation pipeline and data movement pipeline from the L0C buffer to the UB in some hardware architectures.

PIPE_M

Matrix computation pipeline.

PIPE_MTE1

Data movement pipeline from the L1 buffer to the L0A buffer and from the L1 buffer to the L0B buffer.

PIPE_MTE2

Data movement pipeline from the GM to the L1 buffer, from the GM to the UB, and others.

PIPE_MTE3

Data movement pipeline from the UB to the GM and others.

PIPE_FIX

Data movement pipeline from the L0C buffer to the GM, from the L0C buffer to the L1 buffer, and others.

When calling C APIs for movement or computation to write operators, you need to insert corresponding synchronization events based on the data dependencies between pipelines. C APIs provide two types of synchronization control APIs, with varying levels of granularity, to help you accurately adapt to the hardware architecture and explore the performance limits of heterogeneous computation.

First type: synchronization APIs consistent with static tensor programming. These APIs are managed through the asc_sync_notify/asc_sync_wait API for fine-grained control. You need to manually manage the event types and event IDs, and consider forward synchronization (intra-loop dependency) and backward synchronization (inter-loop dependency). This type of API is recommended for scenarios requiring ultimate performance. The sample is as follows:

1
2
3
4
5
6
7
8
// This code snippet is used only to describe the relationship between data movement, vector computation, and synchronization operations. For details about the complete parameters and context of each API, see the programming example below.
asc_copy_gm2ub(); // GM -> UB movement pipeline
asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
asc_add(); // Vector computation pipeline
asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
asc_copy_ub2gm(); // UB -> GM movement pipeline

Second 2: synchronization APIs that is unaware of the pipeline type. Add the asc_sync API after the instruction of the corresponding pipeline type. When using this type of synchronization API, you do not need to consider the instruction pipeline type. The APIs automatically manage the synchronization of all instruction pipelines, simplifying the synchronization instructions. This type of API can be used in performance-insensitive scenarios. The sample is as follows:

1
2
3
4
5
6
// This code snippet is used only to describe the relationship between data movement, vector computation, and synchronization operations. For details about the complete parameters and context of each API, see the programming example below.
asc_copy_gm2ub();// GM -> UB movement pipeline
asc_sync(); // Full synchronization, without the need of considering the subsequent instruction pipeline.
asc_add(); // Vector computation pipeline
asc_sync(); // Full synchronization, without the need of considering the subsequent instruction pipeline.
asc_copy_ub2gm(); // UB -> GM movement pipeline

In addition, C APIs provide a group of movement and computation APIs that include synchronization capabilities. You do not need to manually manage synchronization in explicit mode, as synchronization operations are hidden in the corresponding APIs. This type of API is recommended in performance-insensitive scenarios. The sample is as follows:

1
2
3
4
// This code snippet is used only to describe the relationship between data movement, vector computation, and synchronization operations. For details about the complete parameters and context of each API, see the programming example below.
asc_copy_gm2ub_sync(); // The GM -> UB movement pipeline also includes synchronization with any subsequent instruction pipeline.
asc_add_sync(); // The vector computation pipeline also includes synchronization with any subsequent instruction pipeline.
asc_copy_ub2gm_sync(); // The UB -> GM movement pipeline also includes synchronization with any subsequent instruction pipeline.

Programming Sample

The following is a complete example of memory management and fine-grained synchronization:

 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
49
50
51
52
53
54
55
56
57
#include <cstdint>
#include "c_api/asc_simd.h"

constexpr uint32_t C_API_ONE_BLOCK_SIZE = 32;
constexpr uint32_t C_API_ONE_REPEAT_BYTE_SIZE = 256;
constexpr uint32_t C_API_TOTAL_LENGTH = 16384;
constexpr uint32_t C_API_TILE_NUM = 8;
constexpr uint32_t C_API_TILE_LENGTH = 256;

__vector__ __global__ __aicore__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)
{
    asc_init();

    uint32_t blockLength = C_API_TOTAL_LENGTH / asc_get_block_num();
    uint32_t tileLength = blockLength / C_API_TILE_NUM;

    __gm__ float* xGm = x + asc_get_block_idx() * blockLength;
    __gm__ float* yGm = y + asc_get_block_idx() * blockLength;
    __gm__ float* zGm = z + asc_get_block_idx() * blockLength;

    __ubuf__ float xLocal[C_API_TILE_LENGTH];
    __ubuf__ float yLocal[C_API_TILE_LENGTH];
    __ubuf__ float zLocal[C_API_TILE_LENGTH];

    uint16_t burst_len = tileLength;
    for (uint32_t i = 0; i < C_API_TILE_NUM; i++) {
        if (i != 0) {
            asc_sync_wait(PIPE_V, PIPE_MTE2, EVENT_ID0);
        }

        burst_len = tileLength * sizeof(float) / C_API_ONE_BLOCK_SIZE;
        asc_copy_gm2ub(xLocal, xGm + i * tileLength, 0, 1, burst_len, 0, 0);    
        asc_copy_gm2ub(yLocal, yGm + i * tileLength, 0, 1, burst_len, 0, 0);

        asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
        asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);

        if (i != 0) {
            asc_sync_wait(PIPE_MTE3, PIPE_V, EVENT_ID0);
        }

        asc_add(zLocal, xLocal, yLocal, tileLength * sizeof(float) / C_API_ONE_REPEAT_BYTE_SIZE, 1, 1, 1, 8, 8, 8);

        if (i != (C_API_TILE_NUM-1)) {
            asc_sync_notify(PIPE_V, PIPE_MTE2, EVENT_ID0);
        }

        asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
        asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);

        asc_copy_ub2gm(zGm + i * tileLength, zLocal, 0, 1, burst_len, 0, 0);

        if (i != (C_API_TILE_NUM-1)) {
            asc_sync_notify(PIPE_MTE3, PIPE_V, EVENT_ID0);
        }
    }
}

The following is a complete example of memory management and synchronization management that is unaware of the pipeline type:

 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
#include <cstdint>
#include "c_api/asc_simd.h"

constexpr uint32_t TILE_LENGTH = 2048;
constexpr uint32_t NUM_BLOCKS = 8;

__vector__ __global__ __aicore__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)
{
    asc_init();

    uint32_t blockLength = NUM_BLOCKS * TILE_LENGTH / asc_get_block_num();

    __gm__ float* xGm = x + asc_get_block_idx() * blockLength;
    __gm__ float* yGm = y + asc_get_block_idx() * blockLength;
    __gm__ float* zGm = z + asc_get_block_idx() * blockLength;

    __ubuf__ float xLocal[TILE_LENGTH];
    __ubuf__ float yLocal[TILE_LENGTH];
    __ubuf__ float zLocal[TILE_LENGTH];

    asc_copy_gm2ub((__ubuf__ void*)xLocal, (__gm__ void*)xGm, blockLength * sizeof(float));
    asc_copy_gm2ub((__ubuf__ void*)yLocal, (__gm__ void*)yGm, blockLength * sizeof(float));
    asc_sync();

    asc_add(zLocal, xLocal, yLocal, blockLength);
    asc_sync();

    asc_copy_ub2gm((__gm__ void*)zGm, (__ubuf__ void*)zLocal, blockLength * sizeof(float));
    asc_sync();
}

The following is a complete example of memory management and using the API with the synchronization capability:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
#include <cstdint>
#include "c_api/asc_simd.h"

constexpr uint32_t TILE_LENGTH = 2048;
constexpr uint32_t NUM_BLOCKS = 8;

__vector__ __global__ __aicore__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)
{
    asc_init();

    __ubuf__ float xLocal[TILE_LENGTH];
    __ubuf__ float yLocal[TILE_LENGTH];
    __ubuf__ float zLocal[TILE_LENGTH];

    uint32_t blockLength = TILE_LENGTH * NUM_BLOCKS / asc_get_block_num();

    asc_copy_gm2ub_sync((__ubuf__ void*)xLocal, (__gm__ void*)(x + asc_get_block_idx() * blockLength), blockLength * sizeof(float));
    asc_copy_gm2ub_sync((__ubuf__ void*)yLocal, (__gm__ void*)(y + asc_get_block_idx() * blockLength), blockLength * sizeof(float));

    asc_add_sync(zLocal, xLocal, yLocal, blockLength);

    asc_copy_ub2gm_sync((__gm__ void*)(z + asc_get_block_idx() * blockLength), (__ubuf__ void*)zLocal, blockLength * sizeof(float));
}

The following is a complete example of memory management, Reg vector computation, and fine-grained synchronization:

 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
49
50
51
52
53
54
55
56
57
58
59
#include <cstdint>
#include "c_api/asc_simd.h"

onstexpr uint32_t TILE_LENGTH = 2048;
constexpr uint32_t NUM_BLOCKS = 8;
constexpr uint32_t BLK_NUM = 1;
constexpr uint32_t MASK = 32;

__simd_vf__ inline void AddVF(uint16_t rep, uint16_t one_rep_size, uint32_t blockLength, __ubuf__ float* xLocal, __ubuf__ float* yLocal, __ubuf__ float* zLocal)
{
    vector_bool vmask;
    vector_float reg_src0;
    vector_float reg_src1;
    vector_float reg_dst;
    uint32_t remaining = blockLength;
    for (uint16_t i = 0; i < rep; ++i) {
        vmask = asc_update_mask_b32(remaining);
        asc_loadalign(reg_src0, xLocal + i * one_rep_size);
        asc_loadalign(reg_src1, yLocal + i * one_rep_size);    
        asc_add(reg_dst, reg_src0, reg_src1, vmask);
        asc_storealign(zLocal + i * one_rep_size, reg_dst, vmask);
    }
}

__vector__ __global__ __aicore__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)
{
    asc_init();

    uint32_t blockLength = TILE_LENGTH * NUM_BLOCKS / asc_get_block_num();

    __gm__ float* xGm = x + get_block_idx() * blockLength;
    __gm__ float* yGm = y + get_block_idx() * blockLength;
    __gm__ float* zGm = z + get_block_idx() * blockLength;

    __ubuf__ float xLocal[TILE_LENGTH];
    __ubuf__ float yLocal[TILE_LENGTH];
    __ubuf__ float zLocal[TILE_LENGTH];

    const uint8_t cacheMode0 = static_cast<uint8_t>(((uint64_t)xGm) >> 60);
    const uint8_t cacheMode1 = static_cast<uint8_t>(((uint64_t)yGm) >> 60);
    const uint8_t cacheMode2 = static_cast<uint8_t>(((uint64_t)zGm) >> 60);
    uint32_t burstLength = blockLength * 32;
    uint64_t srcStride = burstLength;
    uint32_t dstStride = (burstLength + 31) / 32 * 32;

    asc_copy_gm2ub_align((__ubuf__ float*)xLocal, xGm, BLK_NUM, burstLength, 0, 0, true, cacheMode0, srcStride, dstStride);
    asc_copy_gm2ub_align((__ubuf__ float*)yLocal, yGm, BLK_NUM, burstLength, 0, 0, true, cacheMode1, srcStride, dstStride);
    asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
    asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);

    uint16_t mask_bit_size = 256;
    uint16_t one_rep_size = mask_bit_size / sizeof(float);
    uint16_t rep = (blockLength + one_rep_size - 1) / one_rep_size;
    asc_vf_call<AddVF>(rep, one_rep_size, blockLength, (__ubuf__ float*)xLocal, (__ubuf__ float*)yLocal, (__ubuf__ float*)zLocal );
    asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
    asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);

    asc_copy_ub2gm_align(zGm, (__ubuf__ float*)zLocal, BLK_NUM, burstLength, cacheMode2, srcStride, dstStride);
}
Parent topic: SIMD Programming
This site uses cookies to improve your browsing experience. By continuing to browse this site, you agree to our use of cookies.View Details