Ascend C 算子开发学习笔记
收藏回复举报
Ascend C 算子开发学习笔记
新人帖
发表于2024-07-11 12:22:04
0 查看

扶稳坐好,三分钟带你了解Ascend C

——算子开发与模型调优部署的利器

目录:

  • 基本知识
  • Asend C
  • 算子基本概念
  • 自定义算子
  • 核函数
  • Kernel侧算子
  • Host侧算子
  • 算子优化
  • 基本知识

昇腾AI处理器:

采用自研达芬奇架构+高集成的SoC设计

实现性能、功耗、成本的三维均衡

cke_24997.png

昇腾AI异构计算架构CANN

NPU负责调度;CPU负责同步信息。

特点/优势:

支持融合算子库,支持cpu生态向npu高效迁移,Ascend C支持算子极简开发,图编译加速技术(串行->并行;算子融合;计算图下沉:计算图分割子图,交付相应的芯片)

  • Asend C

概述:

Ascend C开发的算子运行在AI Core上,AI Core是昇腾NPU硬件平台的计算核心,NPU内部有多个AI Core。多个AI Core共享相同的指令代码,每个核上的运行实例唯一的区别是block_idx不同,开发者只需要关注单核上的处理程序——核函数

Ascend C提供流水线式的编程范式,基于编程范式可以快速搭建算子实现的代码框架,实现流水并行。Ascend C把算子核内的处理程序,分成多个流水任务:“搬入、计算、搬出”,通过队列完成任务间通信和同步并通过统一的内存管理模块管理任务间通信内存。

关键部分:

AI Core(升腾AI处理器计算核心):

计算单元:标量、向量、矩阵。

存储系统:多层级

控制单元:提供指令控制,并行控制时间同步。

cke_7409.png

  • 算子基本概念

什么是算子:

算子名称:一个算子的表示,用于定位算子出现的问题

算子类型:如卷积层(Convolution Layer)是一个算子;全连接层(Full-connected Layer FC layer)中的权值求和过程,是一个算子

数据容器(张量Tensor):存储输入/输出数据的容器。名称、形状、数据类型、数据排布格式

cke_464954.png

数据排布格式:怎么存决定怎么取,不同存储方式实现不同功能。

在昇腾AI处理器中,为了提高数据的访问效率,张量数据采用NC1HWC0的五维格式。其中C0与微架构强相关,等于AI Core中矩阵计算单元的大小,这部分数据需要连续存储;C1是将C维度按照C0进行拆分后的数目,即C1=C/C0。如果不整除,最后一份数据需要补齐以对齐C0。

  • 自定义算子:

优势:

基于C/C++编程

屏蔽硬件差异

多层级API封装

孪生调试:CPU可以模拟NPU实现

Host:服务器/主机。Host memary:主机内存。

Device模块:AI加速卡。Device memary:卡内内存。

Local memary核内内存。Globle memary核外内存。

  • 核函数

调用核函数:

kernel_name<<<blockDim(几个核上进行),12ctrl(保留参数),stream(任务队列)>>>(argument list)

blockDim:几个核上进行

12ctrl:保留参数

stream任务队列

  • Kernel侧算子

Init():数据偏移  

process():执行命令

  • Host侧算子

Tiling实现:计算数据切分的参数。

Shape推导:算子构图阶段可以提前静态申请内存。

算子原型注册。

MatMul  计算公式:C=A*B+Bias  、矩阵切分操作、原理

具体切分原理可见:MindSpore

  • 算子优化

  考察两个指标:计算&搬运

  算子构成:数学公式、计算机算法、编程

~Tiling关键是如何切分

~Scalar优化:减少标量运算,如用Host中计算,Tiling搬运到Device

~0API指令:十分灵活,但是参数多

~Shape 32B对齐.

本帖最后由 匿名用户2024/07/11 18:03:50 编辑

我要发帖子