昇腾社区首页
中文
注册
开发者
下载
支持
积分兑换 NEW

基于语言扩展层C API编程

基于语言扩展层C API编程时，通过提供纯C风格的接口，符合C语言算子开发习惯，提供与业界类似编程体验。本节主要介绍C API编程范式，通过内存管理、同步控制、计算及搬运接口相关的介绍，使开发者更好地理解和使用C API进行编程。

内存管理

C API通过C风格的地址限定符描述不同层级内存，并且可以通过指针直接操作内存地址，从而精准控制数据存放位置。不同存储单元的地址限定符介绍如下：

表1 不同存储单元的地址限定符

存储单元

地址限定符

描述

Global Memory

__gm__

表示被修饰的变量位于Global Memory地址空间。

Unified Buffer

__ubuf__

表示被修饰的变量位于Unified Buffer地址空间。

L1 Buffer

__cbuf__

表示被修饰的变量位于L1 Buffer地址空间。

L0A Buffer

__ca__

表示被修饰的变量位于L0A Buffer地址空间。

L0B Buffer

__cb__

表示被修饰的变量位于L0B Buffer地址空间。

L0C Buffer

__cc__

表示被修饰的变量位于L0C Buffer地址空间。

地址空间限定符可以在数组或指针变量声明中使用，用于指定对象分配的区域。同一个类型上不允许使用多个地址空间限定符。

基于C API编程时，开发者需要自行通过显式的内存管理来控制内存，不同层级的内存申请介绍如下：

  • 全局内存（Global Memory）：一般通过Device侧接口aclrtMalloc接口分配传入，需要增加对应地址限定符使用。
  • 内部存储（包含Unified Buffer、L1 Buffer等）：由用户自行申请空间，通过地址限定符关键字在Kernel内声明。无自动垃圾回收机制，需开发者严格控制生命周期。以申请UB空间为例：
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
// 在数组变量声明中使用地址空间限定符
// total_length 指参与计算的数据长度
constexpr uint64_t total_length = 256;
__ubuf__ float xLocal[ total_length ];
__ubuf__ float yLocal[ total_length ];
__ubuf__ float zLocal[ total_length ];

// 在指针变量声明中使用地址空间限定符
uint64_t offset = 0;                                   // 首先为src0申请内存，从0开始。
__ubuf__ half* src0 = (__ubuf__ half*)asc_get_phy_buf_addr(offset);    // 获取src0的地址，通过__ubuf__关键字指定该地址指向UB内存。

同步控制

NPU内部有不同的计算单元，在计算前往往需要把计算数据搬运到计算单元上。不同计算单元上的计算过程、数据搬运过程可划分为不同的流水线。如下表所示：

表2 指令流水类型和相关说明

流水类型

含义

PIPE_S

标量流水线

PIPE_V

矢量计算流水及部分硬件架构下的L0C Buffer->UB数据搬运流水

PIPE_M

矩阵计算流水

PIPE_MTE1

L1 Buffer ->L0A Buffer、L1 Buffer->L0B Buffer数据搬运流水

PIPE_MTE2

GM->L1 Buffer、GM->UB等数据搬运流水

PIPE_MTE3

UB->GM等数据搬运流水

PIPE_FIX

L0C Buffer->GM、L0C Buffer ->L1等数据搬运流水

在调用C API提供的搬运或者计算类API编写算子时，需要根据流水线之间的数据依赖关系插入对应的同步事件。C API提供了两种不同的同步控制接口，同步控制粒度由浅到深，帮助开发者精准适配硬件架构，挖掘异构计算的性能极限。

第一种：和静态Tensor编程方式一致的同步接口，主要通过asc_sync_notify/asc_sync_wait接口来精细化管理，需要手动管理事件的类型和事件ID，还需要考虑正向同步（循环内依赖）与反向同步（循环间依赖）。极致性能场景推荐使用此方式。使用示例如下：

1
2
3
4
5
6
7
8
// 本片段仅用于说明数据搬运、矢量计算、同步操作间的关系。各接口的完整参数及上下文请参考下文中的编程示例。
asc_copy_gm2ub(); // GM->UB的搬运流水
asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
asc_add(); // 矢量计算流水
asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
asc_copy_ub2gm(); // UB->GM的搬运流水

第二种：不感知流水类型的同步接口，将asc_sync接口添加在对应流水类型的指令后面来实现。使用这类同步接口时，不需要考虑指令流水类型，接口内部会自动管理所有指令流水的同步，简化同步指令。性能不敏感场景下，可以使用此方式。使用示例如下：

1
2
3
4
5
6
// 本片段仅用于说明数据搬运、矢量计算、同步操作间的关系。各接口的完整参数及上下文请参考下文中的编程示例。
asc_copy_gm2ub();// GM->UB的搬运流水
asc_sync(); // 全同步 无需考虑后面的指令流水
asc_add(); // 矢量计算流水
asc_sync(); // 全同步 无需考虑后面的指令流水
asc_copy_ub2gm(); // UB->GM的搬运流水

另外，C API还提供了一组包含同步能力的搬运及计算接口，开发者无需显式手动管理同步，同步操作隐藏在相应的接口中。性能不敏感场景下，推荐使用此方式。使用示例如下：

1
2
3
4
// 本片段仅用于说明数据搬运、矢量计算、同步操作间的关系。各接口的完整参数及上下文请参考下文中的编程示例。
asc_copy_gm2ub_sync(); // GM->UB的搬运流水同时包含了和后面的任意指令流水的同步
asc_add_sync(); // 矢量计算流水同时包含了和后面的任意指令流水的同步
asc_copy_ub2gm_sync(); // UB->GM的搬运流水同时包含了和后面的任意指令流水的同步

编程示例

内存管理与精细化同步完整示例：

 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);
        }
    }
}

内存管理与不感知流水类型的同步管理完整示例如下：

 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();
}

内存管理与使用带同步能力的接口完整示例如下：

 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));
}
父主题： AI Core SIMD编程