Ascend C算子性能优化实用技巧05——API使用优化

news/2024/11/19 13:57:57/

Ascend C是CANN针对算子开发场景推出的编程语言,原生支持C和C++标准规范,兼具开发效率和运行性能。使用Ascend C,开发者可以基于昇腾AI硬件,高效的实现自定义的创新算法。

目前已经有越来越多的开发者使用Ascend C,我们将通过几期“Ascend C算子性能优化”专题分享,围绕开发者最为关心的算子性能优化环节,介绍Ascend C算子常用的优化技巧,帮助开发者自主构建出更优性能的算子。专题内容将围绕流水优化、搬运优化、内存优化、API使用优化以及Tiling优化等优化技巧,从方案讲解、优化案例、性能对比等多角度展开介绍。前期内容回顾:

  1. 《Ascend C算子性能优化实用技巧01——流水优化》
  2. 《Ascend C算子性能优化实用技巧02——内存优化》
  3. 《Ascend C算子性能优化实用技巧03——搬运优化》
  4. 《Ascend C算子性能优化实用技巧04——Tiling优化》

下面进入第五期内容,介绍API的相关使用技巧:

  1. 在算子kernel类外创建和初始化TPipe对象,触发类内Scalar编译优化
  2. 通过TQueBind接口,实现VECIN和VECOUT内存复用
  3. 通过SetMaskCount接口开启Counter模式,简化mask处理逻辑
  4. Matmul使能AtomicAdd选项,减少Vector计算
  5. 合理使用归约指令,兼顾累加效率和执行性能

在算子kernel类外创建和初始化TPipe对象,触发类内Scalar编译优化

程序在创建类对象时,会分配内存空间,用于存储类中的相关成员变量或函数。当类中变量需要参与计算时,变量值从内存被加载到寄存器,计算完成后,变量从寄存器存储回内存。Scalar常量折叠和常量传播是编译器编译时的优化方式,优化前编译器会判断变量是否只初始化过一次或只赋值过一次,若满足此编译优化的前提条件,变量值将会尽量驻留在寄存器中,从而在后续使用变量时,将减少读取内存的操作,提升运行性能。

TPipe是用来管理全局内存和同步的框架,用户可以调用TPipe的接口,为TQue/TBuf进行内存分配。在编写Ascend C算子过程中,经常用一个类存放计算所需的相关变量,这里称类名为KernelExample。当TPipe对象在KernelExample类的实现中定义并初始化时,TPipe对象的内存空间在整个KernelExample对象的内存空间之中;需要注意的是,创建TPipe对象时,对象初始化会设置全局变量的TPipe指针,这导致KernelExample对象的内存有被外部污染的风险,此时编译器的编译优化将采取保守策略,不会对KernelExample对象中的Scalar变量进行常量折叠和常量传播。因此,在任何场景下,我们都建议将TPipe对象创建于KernelExample类外部,使得TPipe对象的内存空间独立于KernelExample类对象的内存空间,触发编译器对KernelExample类内Scalar的编译优化,减少算子Scalar指令耗时。

【反例】

代码中TPipe对象由KernelExample类内部创建并初始化,影响编译器scalar折叠优化,在npu侧导致scalar无谓增加。

template <typename ComputeT> class KernelExample {public:__aicore__ inline KernelExample() {}__aicore__ inline void Init(...){...pipe.InitBuffer(xxxBuf, BUFFER_NUM, xxxSize);...}private:...TPipe pipe;...};extern "C" __global__ __aicore__ void example_kernel(...){...KernelExample<float> op;op.Init(...);...}

【正例】

改为由kernel入口函数创建TPipe对象,在KernelExample类中保存TPipe指针使用。

template <typename ComputeT> class KernelExample {public:__aicore__ inline KernelExample() {}__aicore__ inline void Init(..., TPipe* pipeIn){...pipe = pipeIn;pipe->InitBuffer(xxxBuf, BUFFER_NUM, xxxSize);...}private:...TPipe* pipe;...};extern "C" __global__ __aicore__ void example_kernel(...){...TPipe pipe;KernelExample<float> op;op.Init(..., &pipe);...}

【性能对比】

图1.aiv_scalar_time优化前后对比

图2.aiv_scalar_ratio优化前后对比

 通过性能数据对比可以看出,scalar time优化明显,平均时间从281us减少到236us,下降17%;平均scalar_time时延占比从21%下降到17%。因此在scalar bound的场景下可以使用此优化措施。

通过TQueBind接口,实现VECIN和VECOUT内存复用

 纯搬运类算子在执行并不涉及实际vector计算,若存在冗余的vector计算,会导致算子整体执行时间变长。这种场景可以使用Ascend C针对纯搬运类算子提供的TQueBind接口,该接口可以将VECIN和VECOUT之间绑定,省略将数据从VECIN拷贝到VECOUT的步骤,从而避免vector的无谓消耗。

 【反例】

此代码片段中存在LocalTensor -> LocalTensor的DataCopy指令,目的是为了保证搬入和搬出之间的流水同步。

template <typename ComputeT> class KernelExample {public:...__aicore__ inline void Process(...){for (int i = 0; i < iLen; ++i) {...  auto iLocal = QueI.AllocTensor<ComputeT>();DataCopy(iLocal, inGm[i * 32], size);QueI.EnQue(iLocal);auto iLocal = QueI.DeQue<ComputeT>();for (int j = 0; j < jLen; ++j) {  ...auto oLocal = QueO.AllocTensor<ComputeT>();DataCopy(oLocal, iLocal, size); // LocalTensor -> LocalTensor的DataCopy指令,以实现数据从VECIN到VECOUT的搬移QueO.EnQue(oLocal);auto oLocal = QueO.DeQue<ComputeT>();DataCopyPad(outGm[j], oLocal, ...);QueO.FreeTensor(oLocal);}QueI.FreeTensor(iLocal);}}private:...  TQue<QuePosition::VECIN, BUFFER_NUM> QueI;TQue<QuePosition::VECOUT, BUFFER_NUM> QueO;...};extern "C" __global__ __aicore__ void example_kernel(...){...op.Process(...);}

 【正例】

将LocalTensor -> LocalTensor的DataCopy指令替换为TQueBind接口,避免将VECIN拷贝到VECOUT的步骤,从而避免了冗余vector计算。

template <typename ComputeT> class KernelExample {public:...__aicore__ inline void Process(...){for (int i = 0; i < iLen; ++i) {...  auto bindLocal = queBind.AllocTensor<ComputeT>();DataCopy(bindLocal, inGm[i * 32], size);queBind.EnQue(bindLocal);auto bindLocal = queBind.DeQue<ComputeT>();for (int j = 0; j < len; ++j) {...DataCopyPad(outGm[j], bindLocal, ...);}queBind.FreeTensor(bindLocal);}}private:...  TQueBind<QuePosition::VECIN, QuePosition::VECOUT, BUFFER_NUM> queBind; // 使用TQueBind替换原来QueI,QueO...};extern "C" __global__ __aicore__ void example_kernel(...){...op.Process(...);}

【性能对比】

图3.aiv_vec_time优化前后对比

 如上图所示,将反例中DataCopy指令替换为TQueBind之后有明显优化。由于省略了数据从VECIN拷贝到VECOUT的步骤,aiv_vec_time几乎缩减为0。

通过SetMaskCount接口开启Counter模式,简化mask处理逻辑

开发者可以通过mask参数进行掩码操作来控制实际参与计算的个数,目前支持Normal和Counter两种模式:

  1. Normal模式:开发者需要自行指定每次迭代内参与计算的元素以及迭代次数,系统默认为Normal模式。

  2. Counter模式:开发者直接传入计算数据量,实际迭代次数由Vector计算单元自动推断。

通过对比我们可以看到,Normal模式虽然具备单次迭代内的mask能力,但使用不如Counter模式方便,Counter不需要开发者感知迭代次数、处理非对齐尾块的操作。

具体来看,Normal模式下,当用户想要指定API计算的总元素个数时,首先需要自行判断是否存在不同的主尾块,主块需要将mask设置为全部元素参与计算,并且计算主块所需迭代次数,然后根据尾块中剩余元素个数重置mask,再进行尾块的运算,这中间涉及大量Scalar计算。Counter模式下,用户不需要计算迭代次数以及判断是否存在尾块,将mask模式设置为Counter模式后,只需要设置mask为{0, 总元素个数},然后调用相应的API,处理逻辑更简便,同时减少代码量和Scalar计算量。

以下反例和正例中的代码均以AddCustom算子为例,修改其中如下Add接口的调用代码,以说明Counter模式的优势。

AscendC::Add(zLocal, xLocal, yLocal, this->tileLength);

【反例】

输入数据类型为half的xLocal, yLocal,数据量均为15000。Normal模式下,每个迭代内参与计算的元素个数最多为256B/sizeof(half)=128个,所以15000次Add计算会被分为:主块计算15000/128=117次迭代,每次迭代128个元素参与计算;尾块计算1次迭代,该迭代15000-117*128=24个元素参与计算。从代码角度,需要计算主块的repeatTimes、尾块元素个数;主块计算时,设置mask值为128,尾块计算时,需要设置mask值为尾块元素个数24;这些过程均涉及Scalar计算。

uint32_t ELE_SIZE = 15000;
AscendC::BinaryRepeatParams binaryParams;uint32_t numPerRepeat = 256 / sizeof(DTYPE_X);  // DTYPE_X为half数据类型
uint32_t mainRepeatTimes = ELE_SIZE / numPerRepeat;
uint32_t tailEleNum = ELE_SIZE % numPerRepeat;AscendC::SetMaskNorm();
AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::NORMAL>(numPerRepeat); // 设置normal模式mask,使每个迭代计算128个数
AscendC::Add<DTYPE_X, false>(zLocal, xLocal, yLocal, AscendC::MASK_PLACEHOLDER, mainRepeatTimes, binaryParams);   // MASK_PLACEHOLDER值为0,此处为mask占位,实际mask值以SetVectorMask设置的为准
if (tailEleNum > 0) {AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::NORMAL>(tailEleNum); // 设置normal模式mask,使每个迭代计算24个数// 偏移tensor的起始地址,在xLocal和yLocal的14976个元素处,进行尾块计算AscendC::Add<DTYPE_X, false>(zLocal[mainRepeatTimes * numPerRepeat], xLocal[mainRepeatTimes * numPerRepeat],  yLocal[mainRepeatTimes * numPerRepeat], AscendC::MASK_PLACEHOLDER, 1, binaryParams);   
}
AscendC::ResetMask();  // 还原mask值

【正例】

输入数据类型为half的xLocal, yLocal,数据量均为15000。Counter模式下,只需要设置mask为所有参与计算的元素个数15000,然后直接调用Add指令,即可完成所有计算,不需要繁琐的主尾块计算,代码较为简练。

当要处理多种15000个元素的矢量计算时,Counter模式的优势更明显,不需要反复修改主块和尾块不同的mask值。

uint32_t ELE_SIZE = 15000;
AscendC::BinaryRepeatParams binaryParams;
AscendC::SetMaskCount();
AscendC::SetVectorMask<DTYPE_X, AscendC::MaskMode::COUNTER>(ELE_SIZE);  // 设置counter模式mask,总共计算15000个数
AscendC::Add<DTYPE_X, false>(zLocal, xLocal, yLocal, AscendC::MASK_PLACEHOLDER, 1, binaryParams);                // MASK_PLACEHOLDER值为0,此处为mask占位,实际mask值以SetVectorMask设置的为准
AscendC::ResetMask();  // 还原mask值

【性能数据】

图4.normal模式 vs counter模式scalar time

图5.normal模式 vs counter模式aiv cycle

图6.normal模式 vs counter模式 total time

从上述三幅性能对比图和示例代码可以看到,使用Counter模式能够大幅度简化代码,易于维护,同时能够降低Scalar和Vector计算耗时,获得性能提升。 

Matmul使能AtomicAdd选项,减少Vector计算

对于Matmul得到的结果矩阵C(m, n),若后续需要和GM上的矩阵D(m, n)进行Add操作,则可以在GetTensorC接口或者IterateAll接口的GM通路上,将enAtomic参数设为1,开启AtomicAdd累加操作,在搬出矩阵C到GM时,矩阵C的结果将直接累加到矩阵D的GM地址上,从而实现与矩阵D的Add操作。

【反例】

将Matmul的结果矩阵C和GM上的矩阵D分别搬到UB上,做完Add操作后,结果再搬出到GM。这样至少要多分配一块UB内存给矩阵D,假设在分离架构的处理器上执行,将多做三次搬运操作(矩阵C从GM搬到UB、矩阵D从GM搬到UB、Add结果从UB搬出到GM)。

 template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE>__aicore__ inline void MatMulKernel(...){...Matmul<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, CFG_MDL> mm;TPipe pipe;REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm);mm.SetTensorA(gm_a);mm.SetTensorB(gm_b);mm.SetBias(gm_bias);mm.IterateAll(local_c);// while (mm.Iterate()) {// mm.GetTensorC(local_c);// }DataCopy(local_d, gm_d, d_size);event_t eventIdMTE2ToV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));SetFlag<HardEvent::MTE2_V>(eventIdMTE2ToV);WaitFlag<HardEvent::MTE2_V>(eventIdMTE2ToV);Add(local_d, local_d, local_c, d_size);DataCopy(gm_d, local_d, d_size);...}extern "C" __global__ __aicore__ void example_kernel(...){...typedef MatmulType<TPosition::GM, CubeFormat::ND, half> aType;  typedef MatmulType<TPosition::GM, CubeFormat::ND, half> bType;  typedef MatmulType<TPosition::GM, CubeFormat::ND, float> cType;  typedef MatmulType<TPosition::GM, CubeFormat::ND, float> biasType;MatMulKernel<aType, bType, cType, biasType)(...);...}

【正例】

计算Matmul结果时,调用IterateAll接口或者GetTensorC接口搬运到矩阵D的GM地址上,同时将接口中enAtomic参数设为1,搬出到GM时,Matmul结果矩阵C会累加到矩阵D上,从而得到两个矩阵Add后的结果。

 template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE>__aicore__ inline void MatMulKernel(...){...Matmul<A_TYPE, B_TYPE, C_TYPE, BIAS_TYPE, CFG_MDL> mm;TPipe pipe;REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm);mm.SetTensorA(gm_a);mm.SetTensorB(gm_b);mm.SetBias(gm_bias);mm.IterateAll(gm_d, 1); // IterateAll接口中的enAtomic设为1// while (mm. Iterate ()) {// mm.GetTensorC(gm_d, 1);     // GetTensorC接口中的enAtomic设为1// }...}extern "C" __global__ __aicore__ void example_kernel(...){...typedef MatmulType<TPosition::GM, CubeFormat::ND, half> aType;  typedef MatmulType<TPosition::GM, CubeFormat::ND, half> bType;  typedef MatmulType<TPosition::GM, CubeFormat::ND, float> cType;  typedef MatmulType<TPosition::GM, CubeFormat::ND, float> biasType;MatMulKernel<aType, bType, cType, biasType)(...);...}

【性能对比】

图7.Matmul使能AtomicAdd选项前后性能对比

以矩阵维度M=64,N=256,K=256,矩阵D为(64, 256)为例,Matmul使能AtomicAdd选项前后的性能对比如上图所示,平均cycle数从开启AtomicAdd选项前的154181变为开启后的135054,性能优化12.4%。因此在这种场景下,使能AtomicAdd选项能获取更优的性能。

合理使用归约指令,兼顾累加效率和执行性能

对于需要将一段连续buffer上的元素全部累加到一个元素上的场景。使用WholeReduceSum的累加效率较高,但是单条指令的执行速度不如BlockReduceSum,因此,为了兼顾累加的效率以及尽量使用执行较快的指令,需要根据不同的shape,通过不同的指令组合达到最佳的性能。

例如在float类型输入,shape大小为256时,使用两次WholeReduceSum或者三次BlockReduceSum都可以得到256个float的累加结果。考虑规约指令之间的组合,一次BlockReduceSum加一次WholeReduceSum也可以完成上述累加效果。由于单条指令,BlockReduceSum执行速度更快,所以性能更优于两次WholeReduceSum,并且由于只使用了两次规约指令,所以性能同样优于三次BlockReduceSum的方案。

【反例】

反例1:

...
static constexpr uint32_t REP_LEN = 256;
TBuf<QuePosition::VECCALC> calcBuf;
pipe.InitBuffer(calcBuf, totalLength * sizeof(float));
AscendC::LocalTensor<float> tempTensor1 = calcBuf.Get<float>();
constexpr uint32_t repCount = REP_LEN / sizeof(float);
const uint32_t repNum0 = (totalLength + repCount - 1) / repCount;
AscendC::SetMaskCount();
AscendC::SetVectorMask<float>(0, totalLength);
AscendC::WholeReduceSum<float, false>(tempTensor1, xLocal, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetVectorMask<float>(0, repNum0);
AscendC::WholeReduceSum<float, false>(zLocal, tempTensor1, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetMaskNorm();
...


反例2

...
constexpr uint32_t c0Count = BLK_LEN / sizeof(DTYPE_X);
const uint32_t blockNum0 = (totalLength + c0Count - 1) / c0Count;
const uint32_t blockNum1 = (blockNum0 + c0Count - 1) / c0Count;
AscendC::SetMaskCount();
AscendC::SetVectorMask<DTYPE_X>(0, totalLength);
AscendC::BlockReduceSum<DTYPE_X, false>(tempTensor1, xLocal, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetVectorMask<DTYPE_X>(0, blockNum0);
AscendC::BlockReduceSum<DTYPE_X, false>(tempTensor1, tempTensor1, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetVectorMask<DTYPE_X>(0, blockNum1);
AscendC::BlockReduceSum<DTYPE_X, false>(zLocal, tempTensor1, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetMaskNorm();
...

【正例】

当输入shape为256,输入类型为float时。采用一次BlockReduceSum加一次WholeReduceSum的组合方式实现将256个float元素求和。完整样例链接请参考ReduceCustom。

...
static constexpr uint32_t BLK_LEN = 32;
TBuf<QuePosition::VECCALC> calcBuf;
pipe.InitBuffer(calcBuf, totalLength * sizeof(float));
AscendC::LocalTensor<float> tempTensor1 = calcBuf.Get<float>();
constexpr uint32_t c0Count = BLK_LEN / sizeof(float);
const uint32_t blockNum0 = (totalLength + c0Count - 1) / c0Count;
AscendC::SetMaskCount();
AscendC::SetVectorMask<float>(0, totalLength);
AscendC::BlockReduceSum<float, false>(tempTensor1, xLocal, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetVectorMask<float>(0, blockNum0);
AscendC::WholeReduceSum<float, false>(zLocal, tempTensor1, AscendC::MASK_PLACEHOLDER, 1,DEFAULT_BLK_STRIDE, DEFAULT_BLK_STRIDE, DEFAULT_REP_STRIDE);
AscendC::PipeBarrier<PIPE_V>();
AscendC::SetMaskNorm();
...

【性能数据】

输入shape为256,数据类型为float,两次WholeReduceSum(反例1)、三次BlockReduceSum(反例2)、一次WholeReduceSum加一次BlockReduceSum(正例),三种累加方式的性能数据(循环100次的时间总和)的性能数据如下:

表1.性能对比

反例1

反例2

正例

13us

13.94us

8.44us

更多学习资源

了解更多Ascend C算子性能优化手段和实践案例,请访问:昇腾社区Ascend C信息专区。


http://www.ppmy.cn/news/1548250.html

相关文章

<项目代码>YOLOv8 瞳孔识别<目标检测>

YOLOv8是一种单阶段&#xff08;one-stage&#xff09;检测算法&#xff0c;它将目标检测问题转化为一个回归问题&#xff0c;能够在一次前向传播过程中同时完成目标的分类和定位任务。相较于两阶段检测算法&#xff08;如Faster R-CNN&#xff09;&#xff0c;YOLOv8具有更高的…

网络安全技术概论知识点

第一章 网络安全基础 知识点 1.网络安全威胁现状、成因&#xff1a; 缺乏相关网络安全法律法规、管理缺失和安全意识不强&#xff1b;政府机构与企业在网络安全方面出发点、思路、侧重点不同、国内外或不同企事业机构及行业等的网络安全标准不统一&#xff1b;网络安全威胁及…

【Linux】软件包管理器yum、编辑器vim

Linux 1.源码安装 和 软件包安装2.软件包管理器yum1.什么是包管理器 3.vim编辑器1.vim常见的三种模式2.vim命令模式&#xff1a;命令集3.vim底行模式&#xff1a;命令集4.使用vim的小技巧5.vim配置 1.源码安装 和 软件包安装 在Linux下安装软件&#xff0c;一个通常的办法是下…

大三学生实习面试经历(1)

最近听了一位学长的建议&#xff0c;不能等一切都准备好再去开始&#xff0c;于是就开始了简历投递&#xff0c;恰好简历过了某小厂的初筛&#xff0c;开启了线上面试&#xff0c;记录了一些问题&#xff1a; &#xff08;通过面试也确实了解到了自己在某些方面确实做的还不够…

centos7 如何卸载和安装达梦数据库实例

1.DM8数据库的卸载和安装 1.1 卸载数据库(卸载和安装部分建议反过来看) 1.1.1 运行uninstall.sh 脚本所在位置为DM8数据库安装所在目录 # 进入DM数据库所在安装目录 cd /dm8 # 运行卸载脚本 ./uninstall.sh 1.1.2 查看安装目录剩下的文件 ll 1.1.3 清空安装目录 #…

Spring Boot汽车资讯:科技与速度的新纪元

摘要 随着信息技术在管理上越来越深入而广泛的应用&#xff0c;管理信息系统的实施在技术上已逐步成熟。本文介绍了汽车资讯网站的开发全过程。通过分析汽车资讯网站管理的不足&#xff0c;创建了一个计算机管理汽车资讯网站的方案。文章介绍了汽车资讯网站的系统分析部分&…

计算机网络——路由选择算法

路由算法 路由的计算都是以子网为单位计算的——找到从原子网到目标子网的路径 链路状态算法 序号——&#xff08;源路由器&#xff0c;序号&#xff09;——如果发现这个序号重复或者老了——就不扩散 先测量——再泛洪获得路由 路由转发情况 若S——>W是21则不更改——…

uni-app页面跳转

2024年8月6日 https://uniapp.dcloud.net.cn/api/router.html#navigateto 非tabBar页面跳转 可用多种方式进行跳转&#xff0c;区别在于对其他页面的处理方式。 uni.navigateTo(OBJECT) 保留当前页面&#xff0c;跳转到应用内的某个非tabBar页面&#xff0c;使uni.navigateBa…