避免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的场景下可以使用此优化措施。
父主题: API使用优化