--- title: 基于语言扩展层C API编程 description: "基于语言扩展层C API编程时，通过提供纯C风格的接口，符合C语言算子开发习惯，提供与业界类似编程体验。本节主要介绍C API编程范式，通过内存管理、同步控制、计算及搬运接口相关的介绍，使开发者更好地理解和使用C API进行编程。" url: https://www.hiascend.com/document/detail/zh/canncommercial/latest/programug/Ascendcopdevg/atlas_ascendc_10_00064.html sourcePath: /source/zh/canncommercial/900/programug/Ascendcopdevg/atlas_ascendc_10_00064.html indexId: 74a76e8008516e03f37b9caa401ec13140cef8d595b4519f36a378e2c85bbfe781 --- # 基于语言扩展层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空间为例： ``` // 在数组变量声明中使用地址空间限定符 // 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，还需要考虑正向同步（循环内依赖）与反向同步（循环间依赖）。极致性能场景推荐使用此方式。使用示例如下： ``` // 本片段仅用于说明数据搬运、矢量计算、同步操作间的关系。各接口的完整参数及上下文请参考下文中的编程示例。 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接口添加在对应流水类型的指令后面来实现。使用这类同步接口时，不需要考虑指令流水类型，接口内部会自动管理所有指令流水的同步，简化同步指令。性能不敏感场景下，可以使用此方式。使用示例如下： ``` // 本片段仅用于说明数据搬运、矢量计算、同步操作间的关系。各接口的完整参数及上下文请参考下文中的编程示例。 asc_copy_gm2ub();// GM->UB的搬运流水 asc_sync(); // 全同步 无需考虑后面的指令流水 asc_add(); // 矢量计算流水 asc_sync(); // 全同步 无需考虑后面的指令流水 asc_copy_ub2gm(); // UB->GM的搬运流水 ``` 另外，C API还提供了一组包含同步能力的搬运及计算接口，开发者无需显式手动管理同步，同步操作隐藏在相应的接口中。性能不敏感场景下，推荐使用此方式。使用示例如下： ``` // 本片段仅用于说明数据搬运、矢量计算、同步操作间的关系。各接口的完整参数及上下文请参考下文中的编程示例。 asc_copy_gm2ub_sync(); // GM->UB的搬运流水同时包含了和后面的任意指令流水的同步 asc_add_sync(); // 矢量计算流水同时包含了和后面的任意指令流水的同步 asc_copy_ub2gm_sync(); // UB->GM的搬运流水同时包含了和后面的任意指令流水的同步 ``` #### 编程示例 内存管理与精细化同步完整示例： ``` #include #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); } } } ``` 内存管理与不感知流水类型的同步管理完整示例如下： ``` #include #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(); } ``` 内存管理与使用带同步能力的接口完整示例如下： ``` #include #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)); } ``` 内存管理、Reg矢量计算与精细化同步完整示例： ``` #include #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(((uint64_t)xGm) >> 60); const uint8_t cacheMode1 = static_cast(((uint64_t)yGm) >> 60); const uint8_t cacheMode2 = static_cast(((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(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); } ```