并行计算架构抽象
指令流、信号流、数据流的基本概念
指令流：指令流是指在计算过程中，指令的执行顺序。在并行计算架构中，如何安排指令流以最大化并行度和资源利用率是一个关键的设计问题。
信号流：信号流涉及到的是在硬件电路中，信号（数据、控制信号等）如何在不同的组件和模块之间传递的方式与路径。在达芬奇架构中，优化信号流是提高效能和响应速度的重要方面。
数据流：数据流关注的是数据在系统中的移动方式，包括数据如何在存储器和处理单元之间传输。高效的数据流设计可以减少等待时间和数据传输的能耗，特别是在数据密集型的AI计算中尤为关键。
AI Core内部并行计算架构抽象
使用Ascend c编程语言开发的算子运行在AI Core上，Al Core是异腾AI处理器中的计算核心个AI处理器内部有多个AI Core，Al Core中包含计算单元、存储单元、搬运单元等核心组件
- 计算单元包括了三种基础计算资源
- Scalar计算单元: 执行地址计算、循环控制等标量计算工作 ，并把向量计算、矩阵计算、数据搬运、同步指令发射给对应单元执行
- Cube计算单元: 负责执行矩阵运算
- Vector计算单元: 负责执行向量运算
- 搬运单元
负责在 Global Memory和Local Memory之间搬运数据
1. MTE2 ( MemoryTransferEngine，数据搬入单元)
2. MTE3 (数据搬出单元) - 存储单元
- 为AI Core的内部存储，统称为Local Memory
- AICore的外部存储称之为Global Memory
三个流
- 异步指令流
Scalar计算单元读取指令序列，并把向量计算、矩阵计算、数据搬运指令发射给对应单元的指令队列向量计算单元、矩阵计算单元、数据搬运单元异步的并行执行接收到的指令 - 同步信号流
指令间可能会存在依赖关系，为了保证不同指令队列间的指令按照正确的逻辑关系执行，scalar计算单元也会给对应单元下发同步指令 - 计算数据流
DMA搬入单元把数据搬运到 Local Memory,Vector/cube计算单元完成数据计算，并把计算结果写回Local Memory，DMA搬出单元把处理好的数据搬运回 Global Memory
SPMD编程模型介绍
什么是SPMD编程
Ascend c算子编程是SPMD的编程，将需要处理的数据拆分并分布在多个计算核心上运行多个AI Core共享相同的指令代码，每个核上的运行实例唯一的区别是block idx不同block的类似于进程，block idx就是标识进程唯一性的进程D，编程中使用函数GetBlockldx()获取ID
异腾AI处理器如何进行SPMD并行计算
好几个batch在AI Core X上分别进行
在进行SPMD并行计算时，异腾AI处理器可以采取以下策略：
-
划分数据集：将大型数据集分割成小块，分配给不同的计算核心。这种数据拆分策略允许并行执行相同的运算任务，但在不同的数据子集上执行，从而提高整体计算效率。
-
自定义操作：通过软件开发套件（SDK）和高层次编程模型（如TensorFlow或PyTorch上的编程接口），开发者可以定义特定的操作，这些操作可以利用异腾AI处理器的并行计算能力，以符合SPMD模型的要求。
-
优化通讯：在并行计算中，不同计算单元之间的通信是一个关键问题。异腾处理器通过高效的通信机制（如专用的互连网络或通信协议）来优化这些通讯操作，减少数据传输时间，确保并行计算的高效率。
-
并行计算优化：异腾处理器通过其硬件设计和软件支持，使得并行计算特别是基于SPMD模型的计算更为高效。例如，通过动态调度和负载平衡，确保每个计算核心都能充分利用其计算能力，避免因个别核心的数据处理延迟而影响整体性能。
核函数编写及调用
什么是核函数
核函数(Kernel Function)是Ascend c算子设备侧的入口。Ascend c允许用户使用核函数这种C/C++函数的语法扩展来管理设备侧的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧和设备侧连接的桥梁
核函数是直接在设备侧执行的代码。在核函数中，需要为在一个核上执行的代码规定要进行的数据访问和计算操作，SPMD编程模型允许核函数调用时，多个核并行地执行同一个计算任务
不同
加了一个__aicore__
如何编写核函数
- 使用函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外，还要为核函数加上额外的函数类型限定符，包含__ global__ 和__ aicore__
使用 __gLobal__函数类型限定符来标识它是一个核函数，可以被 **<<<…>> ** 调用;使用 __ aicore__函数类型限定符来标识该核函数在设备侧AI Core上执行
__global __ __aicore __ void kernel_name(argument list);
| 表1函数类型限定符 |
|---|
| 函数类型限定符 | 执行 | 调用 | 备注 |
|---|---|---|---|
| __global __ | 在设备侧执行 | 由<<<…>>>来调用 | 必须为void返回值类型 |
| __aicore __ | 在设备侧执行 | 仅从设备端调用 | — |
- 使用变量类型限定符
为了方便:指针入参变量统一的类型定义为_gm_uint8_t*
用户可统一使用uint8 t类型的指针，并在使用时转化为实际的指针类型，亦可直接传入实际的指针类型
| 表2变量类型限定符 |
|---|
| 变量类型限定符 | 内存空间 | 意义 |
|---|---|---|
| __gm __ | 驻留在GlobalMemory上 | 表明该指针变量指向Global Memory上某处内存地址 |
-
规则和建议
- 核函数必须具有void返回类型
- 仅支持入参为指针类型或c/C++内置数据类型(Primitive Data Types)，如: half* s0float* s1、int32_tc
- 提供了一个封装的宏GM ADDR来避免过长的函数入参列表
#define GM_ADDR __gm__ uint8_t* __restrict__
如何调用核函数
核函数的调用语句是c/C++函数调用语句的一种扩展
常见的c/C++函数调用方式是如下的形式
function_name(argument list);
核函数使用内核调用符<<<…>>> 这种语法形式，来规定核函数的执行配置:
kernel_name<<<blockDim，l2ctrl，stream>>>(argument list);
注:内核调用符仅可在NPU模式下编译时调用，CPU模式下编译无法识别该符号
- 参数说明
- blockDim，规定了核函数将会在几个核上执行，每个执行该核函数的核会被分配一个逻辑ID，表现为内置变量block idx，编号从0开始，可为不同的逻辑核定义不同的行为，可以在算子实现中使用GetBlockldx()函数来获得
- l2ctrl，保留参数，暂时设置为固定值nullptr
- stream .类型为aclrtStream，stream是一个任务队列，应用程序通过stream来管理任务的并行
调用示例
使用内核调用符<<<…>>>调用核函数:
HelloWorld<<<8，nullptr，stream>>>(fooDevice);
blockDim设置为8，表示在8个核上调用了HelLoworld核函数，每个核都会独立目并行地执行该核函数stream可以通过acLrtCreatestream来创建，它的作用是在当前进程或线程中显式创建一个acLrtstream argument list设置为fooDevice这1个入参
核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机侧强制主机侧程序等待所有核函数执行完毕的API(阻塞应用程序运行，直到指定Stream中的所有任务都完成，同步接口)为aclrtsynchronizestream
aclError aclrtSynchronizestream(aclrtStream stream);
编程API介绍
Ascend C提供了哪些编程API
Ascend 算子采用标准C++语法和一组类库API进行编程
- 计算类API:标量计算API、向量计算AP1矩阵计算AP1，分别实现调用Scalar计算单元、Vector计算单元、Cube计算单元
- 数据搬运API:基于Local Memory数据进行计算，数据需要先从Global Memory搬运至Local Memory，再使用计算接口完成计算，最后从Local Memory搬出至Global Memory。比如DataCopy接口
- 内存管理API:用于分配管理内存，比如A1ocTensor、FreeTensor接口
- 任务同步API: 完成任务间的通信和同步，比如Enoue、Deoue接口。不同的指令异步并行执行，为了保证不同指令队列间的指令按照正确的逻辑关系执行，需要向不同的组件发送同步指令
Ascend CAPI用于计算的基本数据类型都是Tensor: GlobalTensor和LocalTensor
多层级API的特点
4级API定义:API根据用户使用的场景分为4级
3级API，运算符重载，支持+，-，*,/，=，，&，<=实现计算的简单表达，类似dst = src1 + src2
2级连续计算API，类似Add(dst,src1,src2,count)，针对源操作数的连续COUNT个数据进行计算连续写入目的操作数，解决一维tensor的连续count个数据的计算问题
1级slice计算API解决多维数据中的切片计算问题 (开发中)
0级功能丰富计算API，可以完整发挥硬件优势的计算API，该功能可以充分发挥CANN系列芯片的强大功能指令，支持对每个操作数的repeattimes，repeatstride，MASK的操作。调用类似:Add(dst, src1, src2, repeatTimes, repeatParams);
流水编程范式介绍
- 快速开发编程的固定步骤
- 统一代码框架的开发捷径
- 使用者总结出的开发经验
- 面向特定场景的编程思想
- 定制化的方法论开发体验
什么是Ascend C的并行计算编程范式
C编程范式把算子内部的处理程序，分成多个流水任务 (stage)以张量 (Tensor) 为数据载Ascend体，以队列 (Queue) 进行任务之间的通信与同步，以内存管理模块 (Pipe) 管理任务间的通信内存。
抽象编程模型"TPIPE并行计算!
1.针对各代Davinci芯片的复杂数据流，根据实际计算需求，抽象出并行编程范式，简化流水并行2.Ascend C的并行编程范式核心要素
- 一组并行计算任务
- 通过队列实现任务之间的通信和同步
- 程序员自主表达对并行计算任务和资源的调度
3.典型的计算范式
- 基本的矢量编程范式:计算任务分为Copyln,Compute, CopyOut
- 基本的矩阵编程范式:计算任务分为Copyln,Split, Compute, Aggregate, CopyOut
- 复杂的矢量/矩阵编程范式，通过将矢量/矩阵的Out/In 组合在一起的方式来实现复杂计算数据流
什么是流水任务
流水任务(Stage) 指的是单核处理程序中主程序调度的并行任务在核函数内部，可以通过流水任务实现数据的并行处理来提升性能
举例来说，单核处理程序的功能可以被拆分成3个流水任务:
- stage1、stage2、stage3，每个任务专注于完成单一功能;
- 需要处理的数据被切分成n片，使用Progress1~n表示，每个任务需要依次完成n个数据切片的处理。
- Stage间的箭头表达数据间的依赖关系，比如Stage1处理完Progress1之后，Stage2才能对Progress1进行处理
若Progress的n=3，待处理的数据被切分成3片，对于同一片数据，Stage1、Stage2、Stage3之间的处理具有依赖关系，需要串行处理;不同的数据切片，同一时间点，可以有多个流水任务stage在并行处理，由此达到任务并行、提升性能的目的
矢量编程流水任务设计
矢量算子编程范式把算子的实现流程分为3个基本任务: copyIn，Compute，CopyOutCopyIn负责数据搬入操作，compute负责矢量计算操作，copyout负责数据搬出操作
任务间通信和同步
数据通信与同步的管理者
不同的流水任务之间存在数据依赖，需要进行数据传递
Ascend C中使用Queue队列完成任务之间的数据通信和同步，Queue提供了EnQue、DeQue等基础APIQueue队列管理NPU上不同层级的物理内存时，用一种抽象的逻辑位置 (ouePosition)来表达各个级别的存储(storage Scope)，代替了片上物理存储的概念，开发者无需感知硬件架构
矢量编程中Queue类型 (逻辑位置)包括: VECIN、VECOUT
数据的载体
AscendC使用GLobalTensor和LocalTensor作为数据的基本操作单元，它是各种指令API直接调用的对象，也是数据的载体
如何进行矢量编程任务间通信和同步
矢量编程中的逻辑位置(QuePosition): 搬入数据的存放位置: VECIN、搬出数据的存放位置: VECOUT
矢量编程主要分为Copyln、compute、copyOut三个任务:
- Copyln任务中将输入数据从GLobalTensor搬运至LocalTensor后，需要使用EnQue将LocalTensor放入VECIN的Oueue中
- compute任务等待VECIN的Queue中LocalTensor出队之后才可以进行矢量计算，计算完成后使用EnOue将计算结果LocalTensor放入到VECOUT的Oueue中
- CopyOut任务等待VECOUT的Queue中LocalTensor出队，再将其拷贝到GLobalTensor
Stage1: Copyln任务
- 使用DataCopy接口将GLobalTensor拷贝到LocalTensor
- 使用EnQue将LocalTensor放入VECIN的Queue中
Stage2: Compute任务
- 使用DeQue从VECIN中取出LocaLTensor
- 使用Ascend c指令API完成矢量计算:Add
- 使用EnQue将结果LocalTensor放入VECOUT的Queue中
Stage3: CopyOut任务
- 使用DeQue接口从VECOUT的Queue中取出LocaLTensor
- 使用DataCopy接口将LocalTensor拷贝到GLobalTensor
如何管理矢量编程任务中的内存
任务间数据传递使用到的内存统一由内存管理模块Pipe进行管理
Pipe作为片上内存管理者，通过InitBuffer接口对外提供Queue内存初始化功能，开发者可以通过该接口为指定的Queue分配内存
// 使用A1locTensor分配TensorT
Pipe pipe;
TQue<TPosition::VECOUT，2> que;
int num = 4;
int len = 1024;
// InitBuffer分配内存块数为4，每块大小为1024
Bytespipe.InitBuffer(que, num， len);
// A1locTensor分配Tensor长度为1024
BytesLocalTensor<half> tensor1 = que.AlLocTensor();
// 使用FreeTensor释放通过A1locTensor分配的Tensor，注意配对使用
que.FreeTensor<half>(tensor1);
Queue队列内存初始化完成后，需要使用内存时，通过调用ALLocTensor来为LocalTensor分配内存给Tensor，当创建的LocalTensor完成相关计算无需再使用时，再调用FreeTensor来回收LocalTensor的内存
临时变量内存管理
//为TBuf初始化分配内存，分配内存长度为1024字节
TPipe pipe;
TBuf<TPosition::VECIN> calcBuf;
// 模板参数为QuePosition中的VECIN类型
uint32 t byteLen = 1024;
pipe.InitBuffer(calcBuf，byteLen);
// 从calcBuf获取Tensor,Tensor为pipe分配的所有内存大小，为1024字节
LocalTensor<int32_t> tempTensor1 = calcBuf.Get<int32_t>();
// 从calcBuf获取Tensor,Tensor为128个int32_t类型元素的内存大小，为512字节
LocalTensor<int32 t> tempTensor1 = calcBuf,Get<int32 t>(128);
编程过程中使用到的临时变量内存同样通过Pipe进行管理。临时变量可以使用TBuf数据结构来申请指定QuePosition上的存储空间，并使用Get()来将分配到的存储空间分配给新的LocalTensor从TBuf上获取全部长度，或者获取指定长度的LocalTensor
LocalTensor<T> Get<T>();
LocalTensor<T> Get<T>(uint32 t len);
使用TBuf申请的内存空间只能参与计算，无法执行Queue队列的入队出队操作
扫一扫
274
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
