昇腾社区首页
中文
注册

避免TPipe在对象内创建和初始化

【优先级】中

【描述】TPipe是用来管理全局内存和同步的框架,用户可以调用TPipe的接口为TQue/TBuf进行内存分配。scalar常量折叠是在编译时简化常量的一个过程。通过对常量进行折叠,可以简化代码,提升运行性能。参考c++编译,在scalar流水成为性能瓶颈(scalar bound)的场景下可以精简TPipe对象的创建和初始化,减少内存地址的占用,从而减少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的场景下可以使用此优化措施。