写CUDA到底难在哪？ 第1页

总结了一些MindSpore AI框架中写GPU算子的一些经验，供参考

一、在进行CUDA编程之前，需要一些基础知识:
1、并行编程：众多算法都存在其并行的版本，奇偶排序可以将算法复杂度降低到O(n)，双调排序在进一步降低时间复杂度的同时，会带来排序结果的不稳定

2、数学基础：在深度学习领域需要了解数学基础要求并没有那么高，即便如此，在掌握链式求导和矩阵求导之后，也不容易直观的推导出BatchNorm的反向公式。另一方面，如果我们知道Welford算法，对数组一次遍历，即可同时得到mean和variance，这将对性能提升有帮助

3、了解计算机基础原理，也可以帮助我们解决现实中遇到的问题，如：
Softmax计算前，先对分子、分母值求其公约数并做化简，可以有效降低出现溢出的概率；
采用并行规约算法，可以避免浮点数对阶误差，可以提升计算精度；
X86下采用80bit进行double计算，可以解释程序移植到GPU后出现的精度损失的现象

二、当跨过上述问题之后，我们能够写出一个结果正确的CUDA kernel，但是距离“高性能”这一目标还有路要走
1、内存层次结构：Global Memory, Constant Memory, Texture Memory, Share Memory, L1/L2 Cache, Register等。
使用这些内存在得到性能收益的同时，又存在各种的限制，这对程序开发带来了复杂度
以Share Memory为例：

1）相比于Global Memory，Share Memory访问速度快一个数量级
2）Share memory存在容量上的限制：这就需要调整我们的算法，对矩阵切块后计算
3）Share Memory只能Block内部共享：当我们的程序需要实现多个Block之间进行数据同步时，就要另寻他法
4）每个Block过多使用Share Momory时，会降低CUDA Occupancy，这里存在取舍
5）Warp内的Thread访问Share Memory时，程序需要避免Bank Conflict的出现

2、GPU存在若干SM，每个SM包含CUDA Core和Tensor Core等单元，这为我们的程序调度带来了复杂度：
例如，在我们调用CUDA Kernel时，需要配置Grid和Block，什么样的配置是最优的呢？实际上这里没有“万能钥匙”
GPU本身调度存在复杂度：当CUDA kernel启动后，一个Block会被调度到某个SM上，Block中的Thread按照以Warp粒度调度到SM的CUDA Core上。Block之间以及Block内部的Thread之间同时存在并行和串行调度
另一方面，GPU型号众多，输入维度很无法枚举，这给算子算法选择、性能调优带来了海量的搜索空间，对程序泛化性能提出了更高的要求。

三、上面内容仅提及了单个算子性能调优问题，在一个系统中需要考虑的更多问题：

1、在对网络调优时，我们需要考虑算子之间的组合，有时候可能是算子融合，有时候情况更加复杂：

以ReLU算子为例，朴素(也是cuDNN)实现：

ReLU正向：y[i] = x[i] > 0 ? x[i] : 0;

ReLU反向：dx[i] = x[i] > 0 ? dy[i] : 0;

当进一步思考后，可以发现ReLU正反向是IO密集型算子，ReLUGrad访问了x和dy两个矩阵，这是其性能提升的突破口，改进版本如下：

ReLU正向：y[i] = x[i] > 0 ? x[i] : 0;

mask[warp_id] = __ballot_sync(__activemask(), x[i] > 0);

ReLU反向：dx[i] = mask[warp_id][lane_id] > 0 ? dy[i] : 0;

在新版本中，相比于直接将x传递ReLU反向算子，ReLU正向算子给反向算子传递了一个压缩32倍后的mask，这时候稍微牺牲了ReLU正向的性能，降低ReLU反向IO数据量，最终整体性能提升了30%。

2、采用多Stream并行执行，数据拷贝与计算并行，计算与计算并行，可以提升性能

3、Host-Device之间交互，有时需要采用异步编程，从而隐藏Host的延时；有时需要考虑如何对任务进行拆分，充分利用Host和Device各自处理能力

4、在Device内存受限时，也有多种选择：有时候需要将这个算子计算放到Host上；有时候需要将部分内存先放置到Host上，在合适的时间再搬回Device上；有时候将之前的结果丢弃，从而让渡一部分空间，在需要的时候再重新计算

5、再进一步在大规模并行训练中，多机、多卡分布式调优面临更多的问题；同时也需要考虑单点故障等可靠性问题。

更新：MindSpore在GPU上的训练性能优化实践，大家也可以参考