算子性能优化方法介绍¶

提供高性能的计算服务是飞桨的特色之一, 欢迎开发者为飞桨贡献高性能算子, 本文旨在向开发者提供一些快速实现高性能算子的方法。

基本介绍¶

算子性能优化工作的业务范围涵盖前向算子、反向算子、优化器等.
算子性能优化工作的基本目标是获得明显的算子性能提升, 力争达到业界一流的性能水平, 同时保证精度不会下降.
飞桨内算子性能优化主要围绕 GPU 计算开展, 因此需要用户掌握基本的GPU 编程模型.

优化技巧¶

1.通用优化技巧¶

GPU Kernel 直接影响了算子性能, 我们推荐采用以下等通用优化策略提升 GPU Kernel 的性能, 从而削减算子的计算开销.

通用技巧
向量化读写
协线程操作
Warp 级操作
共享内存操作 (注意 Bank Conflicts)

2. 飞桨内置优化技巧¶

我们在飞桨内开发并封装了一些优化技巧, 具体如下表所示, 欢迎使用, 也欢迎在使用过程中提出修改建议.

2.1 线程配置优化 ¶

我们推荐结合 OP 的使用场景设计对于的线程配置策略，如下图所示IndexSample OP常用于处理 2 维数据, 因此使用2 维的线程配置策略相对比 1 维配置策略，性能可提升 20%左右。

优化 GPU Kernel 中的线程配置策略, 涵盖一维、二维、三维线程配置策略, 目前已经在Elementwise, Stack, IndexSample等 OP 中使用.

2.2 Warp 计算优化 ¶

飞桨内对上文中提到的Warp 级操作进行了封装, 提供了简易的调用接口, 开发者可调用接口快速获得 Warp 内或者 Block 内的全部数据的求和、最大值、最小值.

2.3 索引计算优化:¶

当 GPU Kernel 的索引计算中存在除法或取模操作, 将在导致汇编层面计算开销变大, 我们建议采用快速除法优化这部分的计算开销。飞桨内Pooling OP 采用索引优化计算后, 性能提升 1 倍.

飞桨综合了一系列 GPU Kernel 通用性能优化技巧推出了 Kernel Primitive API，提供高性能的 Block 级 IO 运算和 Compute 运算。使用 Kernel Primitive API 进行 Kernel 开发可以更加专注计算逻辑的实现，在保证性能的同时大幅减少代码量，同时实现了算子计算与硬件解耦，详情见官网Kernel Primitive API, 建议参考案例ElementwiseAdd和Reduce 使用。

3. C++模板特性¶

我们也鼓励充分挖掘 C++侧的可用优化点, 如使用#pragma unroll编译阶段加速指令，编译期自动展循环, 加速运行时循环的执行效率.

案例: Elementwise_add OP 采用模板参数加速循环展开, 性能提升约 5%

           struct SameDimsElementwisePrimitiveCaller {
  __device__ inline void operator()(Functor func, ArgsT *args, OutT *result) {
#pragma unroll
    for (int idx = 0; idx < VecSize; ++idx) {
      result[idx] = static_cast<OutT>(Apply(func, args[idx]));
    }
  }
};

          

4. 内置第三方库¶

飞桨内置了 cuBLAS, cuDNN, cuSOLVER, Thrust 等一系列第三方库, 若采用这些第三方等高性能计算库能获得显著的性能收益，也欢迎使用。cuBLAS 使用示例见matmul_kernel_impl.h, cuDNN 的使用示例见conv_kernel.cu, cuSOLVER 使用示例见values_vectors_functor.h, Thrust 使用示例见coalesced_kernel.cu.

算子性能优化 方法介绍¶