网站建设基础书籍,更新标签wordpress,win2008 iis配置网站,网上购物的网站有哪些一、参考资料 【2023 CANN训练营第一季】Ascend C算子开发入门#xff08;中#xff09; 二、重要说明 TIK2编程范式把算子核内的处理程序#xff0c;分成多个流水任务#xff0c;任务之间通过队列#xff08;Queue#xff09;进行通信和同步#xff0c;并通过统一的…一、参考资料 【2023 · CANN训练营第一季】Ascend C算子开发入门中 二、重要说明 TIK2编程范式把算子核内的处理程序分成多个流水任务任务之间通过队列Queue进行通信和同步并通过统一的内存管理模块Pipe管理任务间通信内存。TIK2分别针对Vector、Cube编程设计了不同的流水任务。开发者只需要完成基本任务的代码实现即可底层的指令同步和并行调度由TIK2框架实现开发者无需关注。由于开发高性能Cube算子难度较大当前仅支持用户开发Vector算子。当前TIK2支持的AI处理器型号为昇腾310P AI处理器、昇腾910 AI处理器其他型号暂不支持。当前支持用户使用g等C/C编译器编译在cpu侧执行的TIK2算子并使用gdb单步调试支持用户使用CCEC编译器编译在npu侧执行的TIK2算子实现加速计算暂不支持加载至网络模型中进行整网验证。算子输出的数据类型与输入数据类型相同。输出shape与输入shape相同。 三、相关介绍 1. CANN算子 CANN算子有两种类型TBE算子与AI CPU算子。 AI Core是昇腾AI处理器的计算核心负责执行矩阵、向量、标量计算密集的算子任务在AI Core上执行的算子称为TBETensor Boost Engine算子。AI CPU负责执行不适合跑在AI Core上的算子是AI Core算子的补充主要承担非矩阵类、逻辑比较复杂的分支密集型计算在AI CPU上执行的算子称为AI CPU算子。 1.1 TBE算子 TBETensor Boost Engine张量加速引擎提供了基于TVMTensor Virtual Machine张量虚拟机框架的自定义算子开发能力提供了用户开发自定义算子所需工具。TBE框架给用户提供了两种算子开发方式 DSL与TIK。开发者可以根据需求自由选择两种开发方式的区别如下 DSL Domain-Specific Language 基于特性域语言 DSL接口已高度封装用户仅需要使用DSL接口完成计算过程的表达后续的算子调度、算子优化及编译都可通过已有的接口一键式完成适合初级开发用户。 TIK Tensor Iterator Kernel 张量嵌套内核 开发者可以通过调用TIK提供的API基于Python语言编写自定义算子然后TIK编译器会将其编译为适配昇腾AI处理器应用程序的二进制文件。TIK需要用户手工控制数据搬运和计算流程入门较高但开发方式比较灵活能够充分挖掘硬件能力在性能上有一定的优势。 1.2 AI CPU算子 以下几种场景下可使用AI CPU方式实现自定义算子 不适合跑在AI Core上的算子例如非矩阵类的复杂计算逻辑比较复杂的分支密集型算子等。 例如Dump、profiling等控制算子Queue、Stack等资源状态类算子TopK、Where等检索类算子。 AI Core不支持的算子算子需要某些数据类型但AI Core不支持例如Complex32、Complex64。 某些场景下为了快速打通网络在昇腾AI处理器的执行流程在TBE实现自定义算子较为困难的情况下可通过自定义AI CPU算子进行功能调测提升调测效率。功能调通之后后续性能调测过程中再将AI CPU自定义算子转换为TBE算子实现。 2. TIK TIKTensor Iterator Kernel是一种基于Python语言的动态编程框架呈现为一个Python模块。 开发者可以通过调用TIK提供的API基于Python语言编写自定义算子然后TIK编译器会编译为适配昇腾AI处理器应用程序的二进制文件。 TIK编程模型 使用TIK进行编程的过程如下图所示用户调用TIK API编写算子对应的Python程序后TIK会将其转化为TIK DSL(TIK DSL是一种DSL语言它可以在比CCE更高的抽象层次上定义CCEC程序的行为)经过编译器编译后生成CCEC文件CCEC代码目前对于TIK编程人员无法感知再经过CCE编译器编译后生成可运行在昇腾AI处理器上的应用程序。 3. TIK2 TIK2是一种使用C/C作为前端语言的编程框架开发者可以使用TIK2提供的API编写自定义算子并通过CCEC编译器对自定义算子进行编译生成可运行在昇腾AI处理器上的应用程序。 TIK与TIK2开发方式对比 算子开发方式TIKTIK2语言PythonC/C计算单元AI CoreAI Core编程模型并行化提供串行化编程体系方便编写算子TIK工具自动对计算过程并行化实现高性能。自动内存管理程序员在编写算子的时候不用感知和管理地址编译器会做好内存分配。针对不同的硬件体系结构抽象出统一的并行计算架构屏蔽硬件差异基于抽象的编程架构可以快速开发出高效的算子。调试方式使用TIK调试器进行功能调试可快速定位功能问题。使用gdb工具在CPU侧进行功能调试调试后可无缝移植到AI处理器运行。APIAPI丰富灵活提供高级参数满足高阶用户需求。多层级API封装从简单到灵活兼顾易用与高效。 四、AI Core架构 AI Core是昇腾AI处理器的计算核心可以看成是一个相对简化的现代微处理器的基本架构负责执行矩阵、向量、标量计算密集的算子任务。它包括了三种基础计算资源矩阵计算单元Cube Unit、向量计算单元Vector Unit和标量计算单元Scalar Unit。这三种计算单元各司其职形成了三条独立的执行流水线在系统软件的统一调度下互相配合达到优化的计算效率。AI Core中包含计算单元、存储单元、控制单元、搬运单元。 1. 计算单元 计算单元是AI Core中提供强大算力的核心单元相当于AI Core的主力军主要包括Cube Unit矩阵计算单元、Vector Unit向量计算单元和Scalar Unit标量计算单元完成AI Core中不同类型的数据计算。 计算单元描述CubeCube负责执行矩阵运算。Cube每次执行可以完成一个fp16的1616与1616的矩阵乘例如CAxB如果是int8输入则一次完成16x32与32x16的矩阵乘。其中A来源于L0AB来源于L0BL0C存储矩阵乘的结果和中间结果。VectorVector负责执行向量运算。其算力低于Cube但灵活度高于Cube如支持数学中的求倒数求平方根等。ScalarScalar负责各类型的标量数据运算和程序的流程控制。功能上可以看做一个小CPU完成整个程序的循环控制、分支判断、Cube/Vector等指令的地址和参数计算以及基本的算术运算等。 2. 存储单元 AI Core需要把外部存储中的数据加载到内部存储中才能完成相应的计算。 2.1 外部存储 通常AI Core的外部存储包括L2、HBM、DDR等。 2.2 内部存储 AI Core的内部存储统称为Local Memory主要包括L1 BufferL1缓冲区L0 BufferL0缓冲区Unified Buffer统一缓冲区和Scalar Buffer标量缓冲区。 2.3 存储单元分类 存储单元描述MTEAI Core上有多个MTEMemory Transfer Engine存储转换引擎包括MTE1、MTE2、MTE3。MTE是数据搬运单元负责AI Core内部数据在不同Buffer之间的数据读写管理和格式转换的操作比如填充padding、转置transpose、3D图像转2D矩阵Img2Col等。BIUBIU (Bus Interface Unit总线接口单元)是AI Core的“大门”负责AI Core与总线交互。BIU是AI Core从外部L2缓冲区/双倍速率内存DDR/高速宽带内存HBM读取数据以及往外写数据的出入口负责把AI Core的读写请求转换为总线上的请求并完成协议交互等工作。L1 BufferL1缓冲区通用内部存储是AI Core内比较大的一块数据中转区可暂存AI Core中需要反复使用的一些数据从而减少从总线读写的次数。某些MTE的数据格式转换功能要求源数据必须位于L1 Buffer例如3D图像转2D矩阵Img2Col操作。L0A Buffer / L0B BufferCube指令的输入。L0C BufferCube指令的输出但进行累加计算的时候也是输入的一部分。Unified Buffer统一缓冲区向量和标量计算的输入和输出。Scalar Buffer标量计算的通用缓冲区作为GPR通用寄存器General-Purpose Register不足时的补充。GPR通用寄存器General-Purpose Register标量计算的输入和输出。应用开发工程师不需要具体关注这些寄存器。由系统内部实现封装程序访问Scalar Buffer并执行标量计算的时候系统内部自动实现Scalar Buffer和GPR之间的同步。SPR专用寄存器Special-Purpose RegisterAI Core的一组配置寄存器。通过修改SPR的内容可以修改AI Core的部分计算行为。 2.4 存储单元大小 不同类型的昇腾AI处理器存储单元大小不同用户可通过get_soc_spec接口获取。 2.4.1 函数原型 def get_soc_spec(key) 2.4.2 参数说明 参数名类型说明keystring类型获取硬件信息包含“SOC_VERSION”“AICORE_TYPE”“CORE_NUM”“UB_SIZE”“L2_SIZE”“L1_SIZE”“CUBE_SIZE”“L0A_SIZE”“L0B_SIZE”“L0C_SIZE”“SMASK_SIZE” 2.4.3 返回值 根据输入的key返回对应的值 SOC_VERSION返回标识SOC类型的字符串。AICORE_TYPE返回Core的类型有AiCore或VectorCore两种返回值。CORE_NUM返回核数int类型。UB_SIZE返回UB大小int类型单位Byte。L2_SIZE返回L2大小int类型单位Byte。L1_SIZE返回L1大小int类型单位Byte。CUBE_SIZE返回CUBE大小tuple类型如16,16,16单位为Byte。L0A_SIZE返回L0A大小int类型单位为Byte。L0B_SIZE返回L0B大小int类型单位为Byte。L0C_SIZE返回L0C大小int类型单位为Byte。SMASK_SIZE返回Smask buffer大小int类型单位为Byte。 2.4.4 示例代码 实际调用时将变量soc_version的值修改为实际的昇腾AI处理器型号。 import tbe soc_versionxxx tbe.common.platform.set_current_compile_soc_info(soc_version) tbe.common.platform.get_soc_spec(CORE_NUM)2.5 指令与存储访问关系 上图的存储单元是软件层面概念其中 Scalar Buffer对应硬件存储单元Scalar Buffer。Unified Buffer对应硬件存储单元Unified Buffer。L1 Buffer对应硬件存储单元L1 Buffer。L1Out Buffer为从L0C上抽象出来的存储Cube计算输出数据的存储单元 2.6 QuePosition与硬件存储单元映射关系 QuePosition硬件存储单元GMGlobal MemoryA1L1 BufferA2L0A BufferB1L1 BufferB2L0B BufferCO1L0C BufferCO2Unified Buffer 2.7 硬件存储单元对齐 不同scope的对齐要求如下表所示 scope对齐要求Unified Buffer昇腾310 AI处理器要求32Byte对齐昇腾910 AI处理器要求32Byte对齐昇腾310P AI处理器AI Core要求32Byte对齐昇腾310P AI处理器Vector Core要求32Byte对齐L1 Buffer512Byte对齐L1OUT Bufferhalf类型数据要求512Byte对齐float/int32_t/uint32_t类型数据要求1024Byte对齐Global Memory暂无对齐要求 3. 控制单元 控制单元为整个计算过程提供了指令控制相当于AI Core的司令部负责整个AI Core的运行。系统控制模块System Control负责指挥和协调AI Core的整体运行模式配置参数和实现功耗控制等。当指令通过指令发射模块Instruction Dispatch顺次发射出去后根据指令的不同类型将会分别被发送到矩阵运算队列Cube Queue、向量运算队列Vector Queue和存储转换队列MTE Queue。指令执行过程中可以提前预取后续指令并一次读入多条指令进入缓存提升指令执行效率。多条指令从系统内存通过总线接口BIU进入到AI Core的指令缓存模块Instruction Cache中等待后续硬件快速自动解码或运算。指令被解码后便会被导入标量指令处理队列Scalar PSQ中实现地址解码与运算控制。 3.1 控制单元分类 AI Core包含的控制单元如下表所示。 控制单元描述系统控制模块System Control外部的Task Scheduler控制和初始化AI Core的配置接口 配置PC、Para_base、BlockID等信息具体功能包括Block执行控制、Block执行完之后中断和状态申报、执行错误状态申报等。指令缓存模块Instruction CacheAI Core内部的指令Cache 具有指令预取功能。标量指令处理队列Scalar PSQScalar指令处理队列。指令发射模块Instruction DispatchCUBE/Vector/MTE指令经过Scalar PSQ处理之后地址、参数等要素都已经配置好之后Instruction Dispatch单元根据指令的类型将CUBE/Vector/MTE指令分别分发到对应的指令队列等待相应的执行单元调度执行。矩阵运算队列Cube QueueCube运算队列。同一个队列里的指令顺序执行不同队列之间可以并行执行。向量运算队列Vector QueueVector运算队列。同一个队列里的指令顺序执行不同队列之间可以并行执行。存储转换队列MTE QueueMTE存储转换队列。同一个队列里的指令顺序执行不同队列之间可以并行执行。事件同步模块Event Sync用于控制不同队列指令也叫做不同指令流水之间的依赖和同步的模块。 3.2 指令队列分类 根据调度分类的不同可以把指令分类加上被译码过程直接解释的Scalar指令缩写为S可以有6种指令分类S、V、M、MTE1、MTE2、MTE3。 队列缩写队列名称备注VVector指令队列用于调度向量指令MMatrix指令队列用于调度Cube指令MTE1存储移动指令队列1用于调度如下内存移动指令L1到L0A/L0B/UB或者用SPR初始化L0A/L0B BufferMTE2存储移动指令队列2用于调度如下内存移动指令L2/HBM/DDR到L1/L0A/L0B/UBMTE3存储移动指令队列3用于调度如下内存移动指令UB到L2/HBM/DDR 除S队列之外不同队列的指令能够乱序执行但是队列内部指令为顺序执行即在满足数据依赖的前提下指令的物理执行顺序不一定与代码的书写顺序一致。 硬件按照下发顺序将不同队列的指令分发到相应的队列上执行昇腾AI处理器提供Barrier、set_flag/wait_flag两种指令保证队列内部以及队列之间按照逻辑关系执行。 Barrier本身是一条指令用于在队列内部约束执行顺序。其作用是保证前序队列中所有数据的读写工作全部完成后序指令才能执行。set_flag/wait_flag为两条指令在set_flag/wait_flag的指令中可以指定一对指令队列的关系表示两个队列之间完成一组“锁”机制其作用方式为 set_flag当前序指令的所有读写操作都完成之后当前指令开始执行并将硬件中的对应标志位设置为1。wait_flag当执行到该指令时如果发现对应标志位为0该队列的后续指令将一直被阻塞如果发现对应标志位为1则将对应标志位设置为0同时后续指令开始执行。 注意TBE封装了这种依赖关系所以应用开发人员不必对Barrier或者Flag进行编程。但应用开发人员仍需要理解这个基本原理才能通过合适的代码调度实现更好的同步关系。基于DSL方式进行算子开发无需关注代码调度DSL提供了自动调度auto_schedule机制。 3.3 AI Core指令调度方式 AI Core采用顺序取指令、并行执行指令的调度方式流水线执行过程如下图所示 指令序列被顺序译码。根据指令的类型有两种可能 如果指令是Scalar指令指令会被直接执行。其他指令指令会被调度到5个独立的指令队列然后再分配到某个空间的执行部件执行。 4. 搬运单元 DMA搬运单元负责在Global Memory和Local Memory之间搬运数据具体来说把数据搬运到Local MemoryVector/Cube计算单元完成数据计算并把计算结果写回Local MemoryDMA搬出单元把处理好的数据搬运回Global Memory。DMA搬运单元包括MTE2Memory Transfer Engine数据搬入单元MTE3数据搬出单元。 五、核函数 核函数是直接在Device设备端执行的代码。在核函数中需要为在一个核上执行的代码规定要进行的数据访问和计算操作当核函数被调用时多个核将并行执行同一个计算任务。 extern C __global__ __aicore__ void add_tik2(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z) { }1. 函数类型限定符 编写核函数 核函数的函数类型限定符包括 __global__ 和 __aicore__其中__global__ 标识核函数__aicore__ 表示核函数在设备端aicore上执行。 函数类型限定符执行调用备注global在设备端执行由…来调用必须有一个void返回值类型aicore在设备端执行仅从设备端调用- 2. 变量类型限定符 指针入参变量统一的类型定义为 __gm__ uint8_t*Init()函数的入参统一设置为uint8_t*类型的指针在后续的使用中需要将其转化为实际的指针类型用户亦可直接传入实际的指针类型。 变量类型限定符内存空间意义gm驻留在Global Memory上表明该指针变量指向Global Memory上某处内存地址 3. 核函数调用符 #ifndef __CCE_KT_TEST__ 表示核函数在NPU侧运行核函数通过核函数调用符 ... 调用。... 仅在NPU侧调用在CPU侧直接调用核函数即可。 #ifndef __CCE_KT_TEST__ // call of kernel function void add_tik2_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z) {add_tik2blockDim, l2ctrl, stream(x, y, z); } #endif4. Init()函数实现 constexpr int32_t TOTAL_LENGTH 8 * 2048; // total length of data constexpr int32_t USE_CORE_NUM 8; // num of core used constexpr int32_t BLOCK_LENGTH TOTAL_LENGTH / USE_CORE_NUM; // length computed of each core constexpr int32_t TILE_NUM 8; // split data into 8 tiles for each core constexpr int32_t BUFFER_NUM 2; // tensor num for each queue constexpr int32_t TILE_LENGTH BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // each tile length is seperated to 2 part, due to double buffer__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z) {//获取核函数的输入输出在Global Memory上的内存偏移地址// get start index for current core, core parallelxGm.SetGlobalBuffer((__gm__ half*)x block_idx * BLOCK_LENGTH);yGm.SetGlobalBuffer((__gm__ half*)y block_idx * BLOCK_LENGTH);zGm.SetGlobalBuffer((__gm__ half*)z block_idx * BLOCK_LENGTH);// 通过Pipe内存管理对象为输入输出Queue分配内存// pipe alloc memory to queue, the unit is Bytespipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(half));pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(half));pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(half)); }解释说明 数据整体长度TOTAL_LENGTH为8* 2048平均分配到8个核上运行每个核上处理的数据大小BLOCK_LENGTH为2048。block_idx为核的逻辑ID(__gm__ half*)x block_idx * BLOCK_LENGTH 即为单核处理程序中x在Global Memory上的内存偏移地址。注意因为Init函数的入参统一设置为 uint8_t*这里需要强转成具体的数据类型 (__gm__ half*)再进行偏移。 4.1 BLOCK // 数据整体长度 // total length of data constexpr int32_t TOTAL_LENGTH 8 * 2048;// 使用多核 // num of core used constexpr int32_t USE_CORE_NUM 8;//每个核处理数据的大小 // length computed of each core constexpr int32_t BLOCK_LENGTH TOTAL_LENGTH / USE_CORE_NUM; block_num block_num默认取值为1即不分核而采用分核并行时其取值上限为65535用户需要保证block_num的值不超过此阈值。 在for_range的原型定义里用户通过设置参数block_num来实现分核并行简单代码示例如下 with tik_instance.for_range( 0, 10, block_num10) as i:for_range循环中的表达式会被作用在10个执行实例上最终10个执行实例会被分配到10个核上并行运行每个核拿到一个执行实例和一个不同的Block ID。如果当前可用的核的数量小于10则执行实例会在当前可用的核上分批调度执行如果当前可用的核的数量大于等于10则会根据实际情况调度执行实际运行的核数可能小于等于10。 一个算子中只能调用一次for_range实现分核即设置block_num 2不允许多次开启多核。 CORE_NUM 用户可以通过get_soc_spec接口获取AI Core的个数。 # 请根据实际昇腾AI处理器型号进行设置 soc_versionxxx # 设置昇腾AI处理器的型号及目标核的类型 tbe.common.platform.set_current_compile_soc_info(soc_version) tbe.common.platform.get_soc_spec(CORE_NUM) # 使用该接口前需要先设置芯片类型为保证负载均衡block_num一般尽量设置为实际核数量的倍数。假设芯片内含32个AI Core假如一个张量的形状为16, 2, 32, 32, 32如果以张量的第一维度最外层进行分核则只能绑定16个核。此时可通过将张量的第一维度和第二维度合并使得最外层的长度变成32以此将任务均摊到32个AI Core上使用尽可能多的核并行处理。需要注意的是顾及后端内存自动分配机制限制用户实施分核并行时必须从最外层开始做维度合并。 4.2 Tiling 对于单核上的处理数据可以进行数据切块Tiling。 // split data into 8 tiles for each core constexpr int32_t TILE_NUM 8;// tensor num for each queue constexpr int32_t BUFFER_NUM 2;// each tile length is seperated to 2 part, due to double buffer constexpr int32_t TILE_LENGTH BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; 5. Process()函数实现 基于矢量编程范式将核函数的实现分为3个基本任务CopyInComputeCopyOut。 __aicore__ inline void Process() {// loop count need to be doubled, due to double bufferconstexpr int32_t loopCount TILE_NUM * BUFFER_NUM;// tiling strategy, pipeline parallelfor (int32_t i 0; i loopCount; i) {CopyIn(i);Compute(i);CopyOut(i);} }核函数内通过数据切块Tiling实现流水线之间的并行。举例来说将单核处理数据分成n份使用progress0processn-1表示处理第1n个数据切片。progress0经过CopyIn Stage之后进入Compute StageCopyIn即可以处理progress1做到了流水线间并行。根据编程范式上面的算法分析将整个计算拆分成三个Stage用户单独编写每个Stage的代码三阶段流程示意图如下 5.1 Stage1CopyIn函数实现。 使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。使用EnQue将LocalTensor放入VecIn的Queue中。 __aicore__ inline void CopyIn(int32_t progress) {// alloc tensor from queue memoryLocalTensorhalf xLocal inQueueX.AllocTensorhalf();LocalTensorhalf yLocal inQueueY.AllocTensorhalf();// copy progress_th tile from global tensor to local tensorDataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);// enque input tensors to VECIN queueinQueueX.EnQue(xLocal);inQueueY.EnQue(yLocal); }5.2 Stage2Compute函数实现。 使用DeQue从VecIn中取出LocalTensor。使用TIK2接口Add完成矢量计算。使用EnQue将计算结果LocalTensor放入到VecOut的Queue中。使用FreeTensor将不再使用的LocalTensor进行回收。 __aicore__ inline void Compute(int32_t progress) {// deque input tensors from VECIN queueLocalTensorhalf xLocal inQueueX.DeQuehalf();LocalTensorhalf yLocal inQueueY.DeQuehalf();LocalTensorhalf zLocal outQueueZ.AllocTensorhalf();// call Add instr for computationAdd(zLocal, xLocal, yLocal, TILE_LENGTH);// enque the output tensor to VECOUT queueoutQueueZ.EnQuehalf(zLocal);// free input tensors for reuseinQueueX.FreeTensor(xLocal);inQueueY.FreeTensor(yLocal); }5.3 Stage3CopyOut函数实现。 使用DeQue接口从VecOut的Queue中取出LocalTensor。使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。使用FreeTensor将不再使用的LocalTensor进行回收。 __aicore__ inline void CopyOut(int32_t progress) {// deque output tensor from VECOUT queueLocalTensorhalf zLocal outQueueZ.DeQuehalf();// copy progress_th tile from local tensor to global tensorDataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);// free output tensor for reuseoutQueueZ.FreeTensor(zLocal); }六、Queue通信和同步 任务间通信和同步 不同的流水任务之间存在数据依赖需要进行数据传递。TIK2中使用Queue队列完成任务之间的数据通信和同步提供EnQue、DeQue等基础API。 1. QuePosition逻辑位置 Queue队列管理NPU上不同层级的物理内存时用一种抽象的逻辑位置 (QuePosition) 来表达各个级别的存储(Storage Scope)代替了片上物理存储的概念开发者无需感知硬件架构达到隐藏芯片架构的目的。Queue类型包括VECIN、VECOUT、A1、A2、B1、B2、CO1、CO2其中VECIN、VECOUT主要用于矢量编程具体说明参见[矢量编程](javascript:A1、A2、B1、B2、CO1、CO2用于矩阵编程具体说明参见[矩阵编程](javascript:。 TIK2使用GLobalTensor 和 LocalTensor 作为数据的基本操作单元它是各种指令API直接调用的对象也是数据的载体。 2. 矢量编程 矢量编程中使用到的逻辑位置QuePosition定义如下 搬入数据的存放位置VECIN搬出数据的存放位置VECOUT。 由流水任务设计可知矢量编程主要分为CopyIn、Compute、CopyOut三个任务。 CopyIn任务中将输入数据从Global内存搬运至Local内存后需要使用EnQue将LocalTensor放入VECIN的Queue中Compute任务等待VECIN的Queue中LocalTensor出队之后才可以完成矢量计算计算完成后使用EnQue将计算结果LocalTensor放入到VECOUT的Queue中CopyOut任务等待VECOUT的Queue中LocalTensor出队再将其拷贝到Global内存。这样 Queue队列就完成了三个任务间的数据通信和同步。 具体流程和流程图如下 Stage1CopyIn任务。 使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。使用EnQue将LocalTensor放入VECIN的Queue中。 Stage2Compute任务。 使用DeQue从VECIN中取出LocalTensor。使用TIK2接口完成矢量计算。使用EnQue将计算结果LocalTensor放入到VECOUT的Queue中。 Stage3CopyOut任务。 使用DeQue接口从VECOUT的Queue中去除LocalTensor。使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。 3. 矩阵编程 由流水任务设计可知矩阵编程主要分为CopyInSplitComputeAggregateCopyOut这5个任务。任务间进行数据传递时会使用到的逻辑位置示意图如下 上图中逻辑位置QuePosition定义如下 搬入数据的存放位置A1用于存放整块A矩阵可类比CPU多级缓存中的二级缓存 搬入数据的存放位置B1用于存放整块B矩阵可类比CPU多级缓存中的二级缓存 搬入数据的存放位置A2用于存放切分后的小块A矩阵可类比CPU多级缓存中的一级缓存 搬入数据的存放位置B2用于存放切分后的小块B矩阵可类比CPU多级缓存中的一级缓存 结果数据的存放位置CO1用于存放小块结果C矩阵可理解为Cube Out 结果数据的存放位置CO2用于存放整块结果C矩阵可理解为Cube Out 搬入数据的存放位置VECIN用于矢量计算是否使用根据实际业务需求 搬出数据的存放位置VECOUT用于矢量计算是否使用根据实际业务需求。 具体任务之间的交互流程和流程图如下。 Stage1CopyIn任务。 使用DataCopy接口将GlobalTensor数据拷贝到LocalTensor。使用EnQue将LocalTensor放入A1/B1的Queue中。 Stage2Split任务。 使用DeQue从A1/B1中取出LocalTensor。使用TIK2接口将LocalTensor从A1/B1中搬运到矩阵计算单元。使用EnQue将计算结果LocalTensor放入到A2/B2的Queue中。 Stage3Compute任务。 使用DeQue从A2/B2中取出LocalTensor。使用TIK2接口完成矩阵计算。使用EnQue将计算结果LocalTensor放入到CO1的Queue中。 Stage4Aggregate任务。 使用DeQue从CO1中取出LocalTensor。使用TIK2接口拷贝结果矩阵到CO2。使用EnQue将计算结果LocalTensor放入到CO2的Queue中。 Stage5CopyOut任务。 使用DeQue接口从CO2的Queue中去除LocalTensor。使用DataCopy接口将LocalTensor拷贝到GlobalTensor上。 4. TQue 4.1 EnQue() 将Tensor/TBufHandle push到队列。 4.2 DeQue() 将TBufHandle/Tensor从队列中取出用于后续处理。 七、Pipe内存管理 通过统一的内存管理模块Pipe对任务间数据传递进行管理。 内存初始化Pipe作为片上内存管理者通过 InitBuffer() 接口对外提供Queue内存初始化功能开发者可以通过该接口为指定的Queue分配内存。分配内存Queue队列内存初始化完成后需要使用内存时通过调用 AllocTensor()来为 LocalTensor分配内存当创建的LocalTensor完成相关计算无需再使用时再调用 FreeTensor() 来回收 LocalTensor 的内存。 编程过程中使用到的临时变量内存同样通过Pipe进行管理。临时变量可以使用TBuf数据结构来申请指定QuePosition上的存储空间。使用TBuf申请的内存空间只能参与计算无法执行Queue队列的入队出队操作。具体的接口使用说明请参考TBuf。 InitBuffer() 为指定的Queue分配内存。 八、Vector矢量编程范式 Vector矢量编程范式把算子的实现流程分为3个基本任务CopyInComputeCopyOut。CopyIn负责搬入操作Compute负责矢量指令计算操作CopyOut负责搬出操作。 九、Cube矩阵编程范式 Cube矩阵编程范式把算子的实现流程分为5个基本任务CopyInSplitComputeAggregateCopyOut。CopyIn负责搬入操作Split负责数据切分操作Compute负责矩阵指令计算操作Aggregate负责数据汇聚操作CopyOut负责搬出操作。 十、术语解析 1. GlobalTensor与LocalTensor TIK2使用GlobalTensor和LocalTensor作为数据的基本操作单元它是各种指令API直接调用的对象也是数据的载体。详见数据结构定义。 采用分核并行时L2/HBM/DDR统称Global Memory对每个核均可见。 1.1 GlobalTensor GlobalTensor 存放全局数据支持QuePosition为GM。 1.2 LocalTensor LocalTensor 存放本地数据支持QuePosition为A1, A2, B1, B2, CO1, CO2, SHM。 2. 数据排布格式format 数据排布格式 NCHW和NHWC 在深度学习领域多维数据通过多维数组存储比如卷积神经网络的特征图Feature Map通常用四维数组保存即4D4D格式解释如下 NBatch数量例如图像的数目。HHeight特征图高度即垂直高度方向的像素个数。WWidth特征图宽度即水平宽度方向的像素个数。CChannels特征图通道例如彩色RGB图像的Channels为3。 由于数据只能线性存储因此这四个维度有对应的顺序。不同深度学习框架会按照不同的顺序存储特征图数据比如Caffe排列顺序为[Batch, Channels, Height, Width]即NCHW。TensorFlow中排列顺序为[Batch, Height, Width, Channels]即NHWC。 以一张格式为RGB的图片为例如上图所示。NCHW中C排列在外层实际存储的是“RRRRRRGGGGGGBBBBBB”即同一通道的所有像素值顺序存储在一起而NHWC中C排列在最内层实际存储的则是“RGBRGBRGBRGBRGBRGB”即多个通道的同一位置的像素值顺序存储在一起。 3. 数据类型dtype Tensor对象的数据类型。 取值范围float16, float32, int8, int16, int32, uint8, uint16, bool等。 4. 形状Shape 张量的形状以(D0, D1, … ,Dn-1)的形式表示D0到Dn是任意的正整数。 如形状(3,4)表示第一维有3个元素第二维有4个元素(3,4)表示一个3行4列的矩阵数组。 张量形状描述1(0,)0维张量也是一个标量[1,2,3](3,)1维张量[[1,2],[3,4]](2, 2)2维张量[[[1,2],[3,4]], [[5,6],[7,8]]](2, 2, 2)3维张量 假设有一些照片每个像素点都由红/绿/蓝3色组成即shape里面3的含义照片的宽和高都是20也就是20*20400个像素总共有4张的照片这就是shape(4, 20, 20, 3)的物理含义。 5. 轴axis 轴是相对shape来说的轴代表张量的shape的下标比如张量a是一个5行6列的二维数组即shape是(5,6)则axis0表示是张量中的第一维即行。axis1表示是张量中的第二维即列。 例如张量数据[[[1,2],[3,4]], [[5,6],[7,8]]]Shape为(2,2,2)则轴0代表第一个维度的数据即[[1,2],[3,4]]与[[5,6],[7,8]]这两个矩阵轴1代表第二个维度的数据即[1,2]、[3,4]、[5,6]、[7,8]这四个数组轴2代表第三个维度的数据即12345678这八个数。 轴axis可以为负数此时表示是倒数第axis个维度。 N维Tensor的轴有0 , 1, 2,……N-1。 6. double buffer机制 执行于AI Core上的指令队列主要包括如下几类即矩阵运算队列Cube Queue、向量运算队列Vector Queue和存储转换队列MTE Queue。不同指令队列间的相互独立性和可并行执行特性是double buffer优化机制的基石。 6.1 Unified Buffer统一缓冲区 一个完整的数据搬运和计算过程MTE2将数据从Global Memory搬运到Unified BufferVector完成计算后将结果写回Unified Buffer最后由MTE3将计算结果搬回Global Memory。Vector所有计算的源数据以及目标数据都要求存储在Unified Buffer中并要求32Byte对齐。Unified Buffer数据搬运与Vector计算过程如下图所示 在此过程中数据搬运与Vector计算串行执行Vector计算单元无可避免存在资源闲置问题。举例而言若MTE2、Vector、MTE3三阶段分别耗时t则Vector的时间利用率仅为1/3等待时间过长Vector利用率严重不足。 6.2 double buffer 为减少Vector等待时间double buffer机制将Unified Buffer一分为二即UB_A、UB_B。如下图所示当Vector对UB_A中数据进行读取和计算时MTE2可将下一份数据搬入UB_B中而当Vector切换到计算UB_B时MTE3将UB_A的计算结果搬出而MTE2则继续将下一份数据搬入UB_A中。由此数据的进出搬运和Vector计算实现并行执行Vector闲置问题得以有效缓解。double buffer机制如下图所示 总体来说double buffer是基于MTE指令队列与Vector指令队列的独立性和可并行性通过将数据搬运与Vector计算并行执行以隐藏数据搬运时间并降低Vector指令的等待时间最终提高Vector单元的利用效率用户可以通过在for_range中设置参数thread_num来实现数据并行简单代码示例如下 with tik_instance.for_range(0, 10, thread_num2) as i:注意事项 多数情况下采用double buffer能有效提升Vector的时间利用率缩减算子执行时间。然而double buffer机制缓解Vector闲置问题并不代表它总能带来整体的性能提升。例如 当数据搬运时间较短而Vector计算时间显著较长时由于数据搬运在整个计算过程中的时间占比较低double buffer机制带来的性能收益会偏小。又如当原始数据较小且Vector可一次性完成所有计算时强行使用double buffer会降低Vector计算资源的利用率最终效果可能适得其反。 因此double buffer的性能收益需综合考虑Vector算力、数据量大小、搬运与计算时间占比等多种因素。