Ascend C算子开发学习笔记
收藏回复举报
Ascend C算子开发学习笔记
新人帖
发表于2024-07-18 16:03:35
0 查看
  • Ascend C算子开发基础知识

Ascend C 算子编程语言,使能算子极简开发,能够匹配用户开发习惯、获得最优执行性能、简化算子开发逻辑、提升算子调试效率

Ascend C 支持 SPMD 数据并行,开发者仅需关注单核算子实现

AI Core 是昇腾 AI 处理器的计算核心,采用华为自研的达芬奇架构

达芬奇架构的主要部分:计算单元(包含三种基础计算单元,即矩阵计算单元、向量计算单元、标量计算单元)、存储系统控制单元

算子对应网络中层或者节点的计算逻辑,在数学中的定义为一个函数空间到函数空间上的映射,即 O:X->X

张量(Tensor)是存储算子输入数据与输出数据的容器

Ascend C 算子是 CANN 针对算子开发场景推出的编程语言,通过多层接口抽象自动并行计算孪生调试等关键技术,极大提高算子开发效率,助力AI开发者低成本完成算子开发和模型调优部署。


核函数是 Ascend C 算子设备侧的入口,编写方式为 __global__ __aicore__ void kernel_name(argument list);(注意:核函数必须具有 void 返回类型,仅支持入参为指针类型或 C/C++ 内置数据类型);对比 CUDA 写法:__global__ void kernel_name(argument list);

核函数的调用(使用内核调用符<<<>>>):kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

外部存储GlobalTensor 用来存放 Global Memory(外部存储)的全局数据

内部存储LocalTensor 用来存放 AI Core 中 Local Memory(内部存储)的数据

Ascend C 中使用 Queue 队列完成任务之间的数据通信和同步,Queue 提供了 EnQue、DeQue 等基础 API

Ascend C 编程范式把算子内部的处理程序,分成多个流水任务(Stage),以张量(Tensor)为数据载体,以队列(Queue)进行任务之间的通信与同步,以内存管理模块(Pipe)管理任务间的通信内存。

临时变量可以使用 TBuf 数据结构来申请指定 QuePosition 中的存储空间

Tiling 算法(或 Tiling 策略)是根据算子中不同输入形状确定搬入基本块大小的相关算法


矢量编程任务间通信与同步

矢量编程中的逻辑位置(QuePosition):搬入数据的存放位置为 VECIN、搬出数据的存放位置为 VECOUT;矢量编程主要分为 CopylnComputeCopyOut 三个任务

Copyln 任务中将输入数据从 GlobalTensor 搬运至 LocalTensor 后,需要使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中

Compute 任务等待 VECIN 的 Queue 中LocalTensor出队之后才可以进行矢量计算,计算完成后使用 EnQue 将计算结果 LoadTensor 放入到 VECOUT 的 Queue 中

CopyOut 任务等待 VECOUT 的 Queue 中 LocalTensor 出队,再将其拷贝到 GlobalTensor

我要发帖子