Reg矢量计算编程【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit简介Reg矢量计算API是面向RegBase架构开发的API用户可以通过该API直接对芯片中涉及Vector计算的寄存器进行操作实现更大的灵活性和更好的性能。Reg矢量计算API与基础API功能相似但与基础API输入和输出数据必须为LocalTensor不同Reg矢量计算API的输入或输出数据均为Reg矢量计算寄存器。对于计算类API其功能是从给定的寄存器获取数据进行计算并将结果保存在给定的寄存器。对于搬运类API其功能是实现UB和寄存器的数据搬运。由此可见Reg矢量计算API相较于基础API将数据搬运和Reg计算过程交给用户自主控制从而实现更大的开发自由度。Regbase编程模型基于寄存器Regbase的编程模型支持用户编写和调用Vector Funtion向量函数。这些函数使用__simd_vf__标记并被发送到硬件中的向量运算单元执行。在simd vf函数内部通过Reg矢量计算API实现计算操作其内存层级与编程架构如图1所示。在SIMD Vector的内存架构中最靠近Vector计算单元的是VF Reg它是SIMD的私有内存包含多种类型的Reg矢量计算寄存器用于存放并行处理的多个数据元素。单核内所有的VF Reg寄存器共享一个本地内存资源UB。SIMD架构不支持从全局内存Global Memory加载数据到Reg矢量计算寄存器先将数据从全局内存GM搬运至Unified Buffer再通过显式的Load/Store指令由Unified Buffer加载到Reg矢量计算寄存器中。图 1SIMD Reg矢量计算内存层级SIMD Reg矢量计算编程架构中通过发出指令到Reg矢量计算执行单元执行单元从Registers读取数据进行计算计算结果写回Registers。DMA搬运单元负责在Registers和Local Memory之间搬运数据。图 2SIMD Reg矢量计算编程架构Regbase和Membase编程调用层级在Membase架构中基础API调用框架API或直接调用编译器BuiltIn API实现功能而高阶API则通过调用基础API来实现功能。在Regbase架构中新增Reg矢量计算API用户在算子实现中可以直接调用该API高阶API和基础API也可以调用该API来实现功能Reg矢量计算API则是直接调用编译器BuiltIn API实现功能。在Regbase架构中中间结果可暂存在寄存器中无需数据搬出到Local Memory的开销在Membase架构中所有操作均基于内存进行这意味着每次计算都需要从Local Memory加载数据计算完成后将结果搬回Local Memory中间计算结果都需要暂存在Local Memory上。在Regbase架构中寄存器容纳的最大数据长度为VLVector Length由于寄存器容量的限制每次只能处理VL长度的数据。因此需要对数据进行切分每次从Local Memory搬运VL长度的数据到寄存器中进行计算计算完成后将结果搬回Local Memory。而在Membase架构中则能够直接处理完整长度的LocalTensor无需进行数据切分从而简化了数据处理流程。Reg矢量计算调用层次核函数使用__global__ __aicore__标识的为核函数是Device侧的入口函数Host侧可以通过...语法进行调用。__aicore__函数使用__aicore__标识该函数在Device侧执行。 核函数内可以调用__aicore__函数。simd vf函数使用__simd_vf__标记能被核函数通过simd vf函数调用。simd vf函数内只能调用__simd_callee__函数和constexpr aicore。__simd_callee__子函数在simd vf函数内可以调用子函数并且这些子函数有可能需要返回值或者通过引用传参这类子函数通过__simd_callee__标识。__simd_callee__函数内只能调用__simd_callee__函数和constexpr aicore函数。具体的调用关系图如下以下为唯一合法函数调用链Regbase编程模型中允许定义simd vf函数并且通过__simd_vf__来进行标记这种设计方案有如下优点__aicore__和__simd_vf__代码隔离清晰编译器可以对编译器BuiltIn API的使用范围是否合法做检测。对函数调用做完善的检查报错比如在__simd_vf__内调用__aicore__函数或者simt函数等错误用法。使用__simd_vf__函数编程用户可以控制某些优化选项如多个simd vf函数融合只针对特定函数生效或针对特定函数关闭某些优化。本示例中在__aicore__函数Compute中调用了VF函数AddVF进行向量加法操作。template typename T __aicore__ inline void Compute() { //申请输出队列并读取输入结果 ... //调用simd vf函数 asc_vf_callAddVFT(dstAddr, src0Addr, src1Addr, count, oneRepeatSize, repeatTimes); //写入结果到输出队列并释放输入队列的内存 ... }Reg矢量计算寄存器Reg矢量计算API操作的基础数据类型介绍如下具体API请参考Reg矢量计算。RegTensor矢量数据寄存器Reg矢量计算基本存储单元用于矢量计算。RegTensor的位宽是VLVector Length可存储VL/sizeof(T)的数据T表示数据类型。MaskReg掩码寄存器用于矢量计算中选择参与计算的元素。MaskReg的位宽是VL/8。UnalignRegForLoad UnalignRegForStore非对齐寄存器作为缓冲区用来优化UB和RegTensor之间的连续非对齐地址访问的开销。在读非对齐地址前UnalignReg应该通过LoadUnAlignPre初始化然后再使用LoadUnAlign。在写非对齐地址时先使用StoreUnAlign再使用StoreUnAlignPost进行后处理。AddrReg地址寄存器用于存储地址偏移量的寄存器。AddrReg通过CreateAddrReg初始化然后在循环之中使用AddrReg存储地址偏移量。AddrReg在每层循环中根据所设置的stride进行自增。本示例中的AddVF函数通过Reg矢量计算API的add接口实现两组数据的相加操作实现高效、灵活的向量计算。通过设置MaskReg掩码寄存器根据实际有效数据长度count生成掩码mask控制参与运算的数据元素的数量。通过LoadAlign/StoreAlign接口实现UB和Reg矢量计算寄存器之间的数据搬运。本示例为连续对齐搬入搬出场景使用到的寄存器类型为RegTensor、MaskReg和AddrReg。templatetypename T __simd_vf__ inline void AddVF(__ubuf__ T* dstAddr, __ubuf__ T* src0Addr, __ubuf__ T* src1Addr, uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes) { AscendC::Reg::RegTensorT srcReg0; AscendC::Reg::RegTensorT srcReg1; AscendC::Reg::RegTensorT dstReg; AscendC::Reg::MaskReg mask; AscendC::Reg::AddrReg aReg; for (uint16_t i 0; i repeatTimes; i) { aReg AscendC::Reg::CreateAddrRegT(i, oneRepeatSize); mask AscendC::Reg::UpdateMaskT(count); AscendC::Reg::LoadAlign(srcReg0, src0Addr, aReg); AscendC::Reg::LoadAlign(srcReg1, src1Addr, aReg); AscendC::Reg::Add(dstReg, srcReg0, srcReg1, mask); AscendC::Reg::StoreAlign(dstAddr, dstReg, aReg, mask); } }本示例为连续非对齐搬入搬出场景使用到的寄存器类型为RegTensor、MaskReg、AddrReg以及UnalignRegForLoad和UnalignRegForStore。template typename T __simd_vf__ inline void LoadUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint32_t oneRepeatSize, uint16_t repeatTimes) { AscendC::Reg::RegTensorT srcReg; AscendC::Reg::UnalignRegForLoad ureg0; AscendC::Reg::UnalignRegForStore ureg1; AscendC::Reg::AddrReg aReg; for (uint16_t i 0; i repeatTimes; i) { aReg AscendC::Reg::CreateAddrRegT(i, oneRepeatSize); AscendC::Reg::LoadUnAlignPre(ureg0, srcAddr, aReg); AscendC::Reg::LoadUnAlign(srcReg, ureg0, srcAddr, aReg, 0); AscendC::Reg::StoreUnAlign(dstAddr, srcReg, ureg1, aReg); } AscendC::Reg::StoreUnAlignPost(dstAddr, ureg1, aReg); }流水线同步控制在SIMD的VF函数的编写中有时候需要将不同的值根据循环写入到同一个地址中或者目标dst和源src是同一个地址这就涉及到不同流水的同步指令。SIMD VF函数内不同流水线之间的同步指令使用LocalMemBar来表示。该同步指令指定src源流水线和dst目的流水线如下图所示目的流水线将等待源流水线上所有指令完成才进行执行。写读场景下当写指令使用的寄存器和读指令使用的寄存器相同时可以触发寄存器保序指令将会按照代码顺序执行不需要插入同步指令而当写指令使用的寄存器和读指令使用的的寄存器不同时如果要确保两条指令顺序执行则需要插入同步指令写写场景同理。函数原型template MemType src, MemType dst __simd_callee__ inline void LocalMemBar()如何使用Reg矢量计算API基于寄存器的编程模型是指每次循环将一个VL长度的数据从从LocalTensor通过数据搬运指令加载到寄存器中进行复杂的数学计算Compute后搬出到LocalTensor中所有的计算逻辑均在寄存器中完成从而减少LocalTensor间的数据搬运大大提升了整体性能具体流程如下所示以AddVF函数为例首先定义三个矢量数据寄存器srcReg0、srcReg1和dstReg以及掩码寄存器mask每次将一个VL长度的数据使用数据搬运函数从src0、src1搬入到数据寄存器srcReg0、srcReg1中地址偏移是src0Addr i * oneRepeatSize、src1Addr i * oneRepeatSize然后调用Add函数将结果存入到dstReg中dstReg srcReg0 srcReg1)mask表示参与Add计算的元素个数最后调用数据搬运函数将结果从dstReg中搬出到dst。Add的原型定义如下template typename T DefaultType, MaskMergeMode mode MaskMergeMode::ZEROING, typename U __simd_callee__ inline void Add(U dstReg, U srcReg0, U srcReg1, MaskReg mask)其中模板参数T表示操作数数据类型MaskMergeMode表示mask未筛选的元素在dst中置零或者保留原值UpdateMask函数用于更新参与计算的mask元素每次循环都会消耗一个VL长度的元素。LoadAlign和StoreAlign函数用于数据的搬入搬出LoadAlign(srcReg0, src0Addr i * oneRepeatSize)表示数据从LocalTensor搬入到srcReg0寄存器起始地址是src0Addr i * oneRepeatSizeStoreAlign(dstAddr i * oneRepeatSize, dstReg, mask)表示将dstReg搬出到LocalTensor目标地址是dstAddr i * oneRepatSize, mask表示有多少元素参与搬出。Reg矢量计算编程示例以Add函数为例宏函数AddVF使用__simd_vf__标记这样的函数也被称为SIMD VF函数。AddVF包含6个参数。dstAddr表示输出数据src0Addr和src1Addr表示输入数据。__ubuf__ 类型表示用于矢量计算的Local MemoryUnified Buffer是LocalTensor实际存储的物理位置。count表示输入数据参与计算的元素个数repeatTimes表示循环次数oneRepeatSize表示每次循环参与的数据量。Add函数首先计算每次能搬入到寄存器中的数据量oneRepeatSize和循环次数repeatTimes然后使用GetPhyAddr获取输入数据和输出数据的UB地址并通过asc_vf_callAddVFT调用AddVF宏函数进行计算。// SIMD函数 template typename T __simd_vf__ inline void AddVF( __ubuf__ T* dstAddr, __ubuf__ T* src0Addr, __ubuf__ T* src1Addr, uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes) { AscendC::Reg::RegTensorT src0Reg; AscendC::Reg::RegTensorT src1Reg; AscendC::Reg::RegTensorT dstReg; AscendC::Reg::MaskReg mask; for (uint16_t i 0; i repeatTimes; i) { mask AscendC::Reg::UpdateMaskT(count); AscendC::Reg::LoadAlign(src0Reg, src0Addr i * oneRepeatSize); AscendC::Reg::LoadAlign(src1Reg, src1Addr i * oneRepeatSize); AscendC::Reg::Add(dstReg, src0Reg, src1Reg, mask); AscendC::Reg::StoreAlign(dstAddr i * oneRepeatSize, dstReg, mask); } } template typename T __aicore__ inline void Compute() { AscendC::LocalTensorT dst outQueueZ.AllocTensorT(); AscendC::LocalTensorT src0 inQueueX.DeQueT(); AscendC::LocalTensorT src1 inQueueY.DeQueT(); constexpr uint32_t oneRepeatSize AscendC::GetVecLen() / sizeof(T); uint32_t count 512; // 向上取整计算循环次数 uint16_t repeatTimes AscendC::CeilDivision(count, oneRepeatSize); __ubuf__ T* dstAddr (__ubuf__ T*)dst.GetPhyAddr(); __ubuf__ T* src0Addr (__ubuf__ T*)src0.GetPhyAddr(); __ubuf__ T* src1Addr (__ubuf__ T*)src1.GetPhyAddr(); asc_vf_callAddVFT(dstAddr, src0Addr, src1Addr, count, oneRepeatSize, repeatTimes); outQueueZ.EnQue(dst); inQueueX.FreeTensor(src0); inQueueY.FreeTensor(src1); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
CANN/asc-devkit：Reg矢量计算编程
Reg矢量计算编程【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit简介Reg矢量计算API是面向RegBase架构开发的API用户可以通过该API直接对芯片中涉及Vector计算的寄存器进行操作实现更大的灵活性和更好的性能。Reg矢量计算API与基础API功能相似但与基础API输入和输出数据必须为LocalTensor不同Reg矢量计算API的输入或输出数据均为Reg矢量计算寄存器。对于计算类API其功能是从给定的寄存器获取数据进行计算并将结果保存在给定的寄存器。对于搬运类API其功能是实现UB和寄存器的数据搬运。由此可见Reg矢量计算API相较于基础API将数据搬运和Reg计算过程交给用户自主控制从而实现更大的开发自由度。Regbase编程模型基于寄存器Regbase的编程模型支持用户编写和调用Vector Funtion向量函数。这些函数使用__simd_vf__标记并被发送到硬件中的向量运算单元执行。在simd vf函数内部通过Reg矢量计算API实现计算操作其内存层级与编程架构如图1所示。在SIMD Vector的内存架构中最靠近Vector计算单元的是VF Reg它是SIMD的私有内存包含多种类型的Reg矢量计算寄存器用于存放并行处理的多个数据元素。单核内所有的VF Reg寄存器共享一个本地内存资源UB。SIMD架构不支持从全局内存Global Memory加载数据到Reg矢量计算寄存器先将数据从全局内存GM搬运至Unified Buffer再通过显式的Load/Store指令由Unified Buffer加载到Reg矢量计算寄存器中。图 1SIMD Reg矢量计算内存层级SIMD Reg矢量计算编程架构中通过发出指令到Reg矢量计算执行单元执行单元从Registers读取数据进行计算计算结果写回Registers。DMA搬运单元负责在Registers和Local Memory之间搬运数据。图 2SIMD Reg矢量计算编程架构Regbase和Membase编程调用层级在Membase架构中基础API调用框架API或直接调用编译器BuiltIn API实现功能而高阶API则通过调用基础API来实现功能。在Regbase架构中新增Reg矢量计算API用户在算子实现中可以直接调用该API高阶API和基础API也可以调用该API来实现功能Reg矢量计算API则是直接调用编译器BuiltIn API实现功能。在Regbase架构中中间结果可暂存在寄存器中无需数据搬出到Local Memory的开销在Membase架构中所有操作均基于内存进行这意味着每次计算都需要从Local Memory加载数据计算完成后将结果搬回Local Memory中间计算结果都需要暂存在Local Memory上。在Regbase架构中寄存器容纳的最大数据长度为VLVector Length由于寄存器容量的限制每次只能处理VL长度的数据。因此需要对数据进行切分每次从Local Memory搬运VL长度的数据到寄存器中进行计算计算完成后将结果搬回Local Memory。而在Membase架构中则能够直接处理完整长度的LocalTensor无需进行数据切分从而简化了数据处理流程。Reg矢量计算调用层次核函数使用__global__ __aicore__标识的为核函数是Device侧的入口函数Host侧可以通过...语法进行调用。__aicore__函数使用__aicore__标识该函数在Device侧执行。 核函数内可以调用__aicore__函数。simd vf函数使用__simd_vf__标记能被核函数通过simd vf函数调用。simd vf函数内只能调用__simd_callee__函数和constexpr aicore。__simd_callee__子函数在simd vf函数内可以调用子函数并且这些子函数有可能需要返回值或者通过引用传参这类子函数通过__simd_callee__标识。__simd_callee__函数内只能调用__simd_callee__函数和constexpr aicore函数。具体的调用关系图如下以下为唯一合法函数调用链Regbase编程模型中允许定义simd vf函数并且通过__simd_vf__来进行标记这种设计方案有如下优点__aicore__和__simd_vf__代码隔离清晰编译器可以对编译器BuiltIn API的使用范围是否合法做检测。对函数调用做完善的检查报错比如在__simd_vf__内调用__aicore__函数或者simt函数等错误用法。使用__simd_vf__函数编程用户可以控制某些优化选项如多个simd vf函数融合只针对特定函数生效或针对特定函数关闭某些优化。本示例中在__aicore__函数Compute中调用了VF函数AddVF进行向量加法操作。template typename T __aicore__ inline void Compute() { //申请输出队列并读取输入结果 ... //调用simd vf函数 asc_vf_callAddVFT(dstAddr, src0Addr, src1Addr, count, oneRepeatSize, repeatTimes); //写入结果到输出队列并释放输入队列的内存 ... }Reg矢量计算寄存器Reg矢量计算API操作的基础数据类型介绍如下具体API请参考Reg矢量计算。RegTensor矢量数据寄存器Reg矢量计算基本存储单元用于矢量计算。RegTensor的位宽是VLVector Length可存储VL/sizeof(T)的数据T表示数据类型。MaskReg掩码寄存器用于矢量计算中选择参与计算的元素。MaskReg的位宽是VL/8。UnalignRegForLoad UnalignRegForStore非对齐寄存器作为缓冲区用来优化UB和RegTensor之间的连续非对齐地址访问的开销。在读非对齐地址前UnalignReg应该通过LoadUnAlignPre初始化然后再使用LoadUnAlign。在写非对齐地址时先使用StoreUnAlign再使用StoreUnAlignPost进行后处理。AddrReg地址寄存器用于存储地址偏移量的寄存器。AddrReg通过CreateAddrReg初始化然后在循环之中使用AddrReg存储地址偏移量。AddrReg在每层循环中根据所设置的stride进行自增。本示例中的AddVF函数通过Reg矢量计算API的add接口实现两组数据的相加操作实现高效、灵活的向量计算。通过设置MaskReg掩码寄存器根据实际有效数据长度count生成掩码mask控制参与运算的数据元素的数量。通过LoadAlign/StoreAlign接口实现UB和Reg矢量计算寄存器之间的数据搬运。本示例为连续对齐搬入搬出场景使用到的寄存器类型为RegTensor、MaskReg和AddrReg。templatetypename T __simd_vf__ inline void AddVF(__ubuf__ T* dstAddr, __ubuf__ T* src0Addr, __ubuf__ T* src1Addr, uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes) { AscendC::Reg::RegTensorT srcReg0; AscendC::Reg::RegTensorT srcReg1; AscendC::Reg::RegTensorT dstReg; AscendC::Reg::MaskReg mask; AscendC::Reg::AddrReg aReg; for (uint16_t i 0; i repeatTimes; i) { aReg AscendC::Reg::CreateAddrRegT(i, oneRepeatSize); mask AscendC::Reg::UpdateMaskT(count); AscendC::Reg::LoadAlign(srcReg0, src0Addr, aReg); AscendC::Reg::LoadAlign(srcReg1, src1Addr, aReg); AscendC::Reg::Add(dstReg, srcReg0, srcReg1, mask); AscendC::Reg::StoreAlign(dstAddr, dstReg, aReg, mask); } }本示例为连续非对齐搬入搬出场景使用到的寄存器类型为RegTensor、MaskReg、AddrReg以及UnalignRegForLoad和UnalignRegForStore。template typename T __simd_vf__ inline void LoadUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint32_t oneRepeatSize, uint16_t repeatTimes) { AscendC::Reg::RegTensorT srcReg; AscendC::Reg::UnalignRegForLoad ureg0; AscendC::Reg::UnalignRegForStore ureg1; AscendC::Reg::AddrReg aReg; for (uint16_t i 0; i repeatTimes; i) { aReg AscendC::Reg::CreateAddrRegT(i, oneRepeatSize); AscendC::Reg::LoadUnAlignPre(ureg0, srcAddr, aReg); AscendC::Reg::LoadUnAlign(srcReg, ureg0, srcAddr, aReg, 0); AscendC::Reg::StoreUnAlign(dstAddr, srcReg, ureg1, aReg); } AscendC::Reg::StoreUnAlignPost(dstAddr, ureg1, aReg); }流水线同步控制在SIMD的VF函数的编写中有时候需要将不同的值根据循环写入到同一个地址中或者目标dst和源src是同一个地址这就涉及到不同流水的同步指令。SIMD VF函数内不同流水线之间的同步指令使用LocalMemBar来表示。该同步指令指定src源流水线和dst目的流水线如下图所示目的流水线将等待源流水线上所有指令完成才进行执行。写读场景下当写指令使用的寄存器和读指令使用的寄存器相同时可以触发寄存器保序指令将会按照代码顺序执行不需要插入同步指令而当写指令使用的寄存器和读指令使用的的寄存器不同时如果要确保两条指令顺序执行则需要插入同步指令写写场景同理。函数原型template MemType src, MemType dst __simd_callee__ inline void LocalMemBar()如何使用Reg矢量计算API基于寄存器的编程模型是指每次循环将一个VL长度的数据从从LocalTensor通过数据搬运指令加载到寄存器中进行复杂的数学计算Compute后搬出到LocalTensor中所有的计算逻辑均在寄存器中完成从而减少LocalTensor间的数据搬运大大提升了整体性能具体流程如下所示以AddVF函数为例首先定义三个矢量数据寄存器srcReg0、srcReg1和dstReg以及掩码寄存器mask每次将一个VL长度的数据使用数据搬运函数从src0、src1搬入到数据寄存器srcReg0、srcReg1中地址偏移是src0Addr i * oneRepeatSize、src1Addr i * oneRepeatSize然后调用Add函数将结果存入到dstReg中dstReg srcReg0 srcReg1)mask表示参与Add计算的元素个数最后调用数据搬运函数将结果从dstReg中搬出到dst。Add的原型定义如下template typename T DefaultType, MaskMergeMode mode MaskMergeMode::ZEROING, typename U __simd_callee__ inline void Add(U dstReg, U srcReg0, U srcReg1, MaskReg mask)其中模板参数T表示操作数数据类型MaskMergeMode表示mask未筛选的元素在dst中置零或者保留原值UpdateMask函数用于更新参与计算的mask元素每次循环都会消耗一个VL长度的元素。LoadAlign和StoreAlign函数用于数据的搬入搬出LoadAlign(srcReg0, src0Addr i * oneRepeatSize)表示数据从LocalTensor搬入到srcReg0寄存器起始地址是src0Addr i * oneRepeatSizeStoreAlign(dstAddr i * oneRepeatSize, dstReg, mask)表示将dstReg搬出到LocalTensor目标地址是dstAddr i * oneRepatSize, mask表示有多少元素参与搬出。Reg矢量计算编程示例以Add函数为例宏函数AddVF使用__simd_vf__标记这样的函数也被称为SIMD VF函数。AddVF包含6个参数。dstAddr表示输出数据src0Addr和src1Addr表示输入数据。__ubuf__ 类型表示用于矢量计算的Local MemoryUnified Buffer是LocalTensor实际存储的物理位置。count表示输入数据参与计算的元素个数repeatTimes表示循环次数oneRepeatSize表示每次循环参与的数据量。Add函数首先计算每次能搬入到寄存器中的数据量oneRepeatSize和循环次数repeatTimes然后使用GetPhyAddr获取输入数据和输出数据的UB地址并通过asc_vf_callAddVFT调用AddVF宏函数进行计算。// SIMD函数 template typename T __simd_vf__ inline void AddVF( __ubuf__ T* dstAddr, __ubuf__ T* src0Addr, __ubuf__ T* src1Addr, uint32_t count, uint32_t oneRepeatSize, uint16_t repeatTimes) { AscendC::Reg::RegTensorT src0Reg; AscendC::Reg::RegTensorT src1Reg; AscendC::Reg::RegTensorT dstReg; AscendC::Reg::MaskReg mask; for (uint16_t i 0; i repeatTimes; i) { mask AscendC::Reg::UpdateMaskT(count); AscendC::Reg::LoadAlign(src0Reg, src0Addr i * oneRepeatSize); AscendC::Reg::LoadAlign(src1Reg, src1Addr i * oneRepeatSize); AscendC::Reg::Add(dstReg, src0Reg, src1Reg, mask); AscendC::Reg::StoreAlign(dstAddr i * oneRepeatSize, dstReg, mask); } } template typename T __aicore__ inline void Compute() { AscendC::LocalTensorT dst outQueueZ.AllocTensorT(); AscendC::LocalTensorT src0 inQueueX.DeQueT(); AscendC::LocalTensorT src1 inQueueY.DeQueT(); constexpr uint32_t oneRepeatSize AscendC::GetVecLen() / sizeof(T); uint32_t count 512; // 向上取整计算循环次数 uint16_t repeatTimes AscendC::CeilDivision(count, oneRepeatSize); __ubuf__ T* dstAddr (__ubuf__ T*)dst.GetPhyAddr(); __ubuf__ T* src0Addr (__ubuf__ T*)src0.GetPhyAddr(); __ubuf__ T* src1Addr (__ubuf__ T*)src1.GetPhyAddr(); asc_vf_callAddVFT(dstAddr, src0Addr, src1Addr, count, oneRepeatSize, repeatTimes); outQueueZ.EnQue(dst); inQueueX.FreeTensor(src0); inQueueY.FreeTensor(src1); }【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
