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;矢量编程主要分为 Copyln、Compute、CopyOut 三个任务
Copyln 任务中将输入数据从 GlobalTensor 搬运至 LocalTensor 后,需要使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中
Compute 任务等待 VECIN 的 Queue 中LocalTensor出队之后才可以进行矢量计算,计算完成后使用 EnQue 将计算结果 LoadTensor 放入到 VECOUT 的 Queue 中
CopyOut 任务等待 VECOUT 的 Queue 中 LocalTensor 出队,再将其拷贝到 GlobalTensor
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;矢量编程主要分为 Copyln、Compute、CopyOut 三个任务
Copyln 任务中将输入数据从 GlobalTensor 搬运至 LocalTensor 后,需要使用 EnQue 将 LocalTensor 放入 VECIN 的 Queue 中
Compute 任务等待 VECIN 的 Queue 中LocalTensor出队之后才可以进行矢量计算,计算完成后使用 EnQue 将计算结果 LoadTensor 放入到 VECOUT 的 Queue 中
CopyOut 任务等待 VECOUT 的 Queue 中 LocalTensor 出队,再将其拷贝到 GlobalTensor