CANN/Ascend C 基于语言扩展层C API编程

基于语言扩展层C API编程【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit基于语言扩展层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_MTE1L1 Buffer -L0A Buffer、L1 Buffer-L0B Buffer数据搬运流水PIPE_MTE2GM-L1 Buffer、GM-UB等数据搬运流水PIPE_MTE3UB-GM等数据搬运流水PIPE_FIXL0C 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 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 len_burst tileLength; for (uint32_t i 0; i C_API_TILE_NUM; i) { if (i ! 0) { asc_sync_wait(PIPE_V, PIPE_MTE2, EVENT_ID0); } len_burst tileLength * sizeof(float) / C_API_ONE_BLOCK_SIZE; asc_copy_gm2ub(xLocal, xGm i * tileLength, 0, 1, len_burst, 0, 0); asc_copy_gm2ub(yLocal, yGm i * tileLength, 0, 1, len_burst, 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, len_burst, 0, 0); if (i ! (C_API_TILE_NUM-1)) { asc_sync_notify(PIPE_MTE3, PIPE_V, EVENT_ID0); } } }内存管理与不感知流水类型的同步管理完整示例如下#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(); }内存管理与使用带同步能力的接口完整示例如下#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)); }内存管理、Reg矢量计算与精细化同步完整示例#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_castuint8_t(((uint64_t)xGm) 60); const uint8_t cacheMode1 static_castuint8_t(((uint64_t)yGm) 60); const uint8_t cacheMode2 static_castuint8_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_callAddVF(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); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考