写CUDA到底难在哪？

最新推荐文章于 2023-11-15 10:05:19 发布

woshicver

最新推荐文章于 2023-11-15 10:05:19 发布

阅读量1k

点赞数 1

文章标签：算法 java 编程语言人工智能深度学习

链接：https://www.zhihu.com/question/437131193

编辑：深度学习与计算机视觉

声明：仅做学术分享，侵删

作者：知乎用户
https://www.zhihu.com/question/437131193/answer/1659455614

其实从我这几年的经验来看，最难的是并行思想。并行思想其中难上加难的东西是数据分组（partitioning），这节是绝对影响性能最厉害地方，以及劝退大量靠算法刷题山来的孩子们。

我们写一个普通程序，一般去考虑的首先是写一个朴素（naive）实现，然后再去看这些地方哪里有可以优化的算法，时间降低时间空间复杂度的要求。而并行计算中，最重要的一点是为数据分组成多个小数据块，每个线程（进程）再去实现SPMD或者SIMD/T。而这个数据分组的方式，存储方法等等直接的影响到你这个并行程序最终的性能。恰恰这一块又是普通没有经过训练的程序员平时完全不考虑的地方。而现代复杂的GPU异构系统导致这个数据分组异常的难。比如要使用一个多GPU的机器，你要把数据分成每个GPU的大组，每个GPU要分成每个Thread的小组，然后为了使用张量（Tensor）核心还要分成张量处理器能用的固定数量的小组，这个过程简直是难上艰难，烧脑子掉头发。

事实上，大部分的并行程序，解决了数据分组问题，其本身的问题就解决了，算法本身的优化反倒是不是那么的重要了（我们其实一般称之为MicroOptimization）

作者：金雪锋
https://www.zhihu.com/question/437131193/answer/1685402769

总结了一些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、再进一步在大规模并行训练中，多机、多卡分布式调优面临更多的问题；同时也需要考虑单点故障等可靠性问题。

作者：Keepin
https://www.zhihu.com/question/437131193/answer/1658645679

cuda目前主要事利用GPU的特性进行加速，使用cuda进行加速有几个重要点需要考虑：

1、并行思想

如果一个算法无法通过并行加速，那么可以考虑放弃用cuda实现，因为这样cuda实现没有太大的效益。（当然还有一种想法就是解放cpu，而不在乎性能）。并行思想也是gpu编程的基础

2、独立显卡的数据传输耗时

对于独显还要考虑如何减少host端（cpu）的内存到device端（gpu）的内存传输耗时（独立显卡由自己的显存，一般通过pcie来传输cpu和gpu，数据量大时候耗时明显）

3、算法特性与架构特性

区分是计算密集型还是IO密集型算法，IO密集型的算法：要考虑GPU的缓存/共享内存如何使用来达到减少带宽传输耗时；还要考虑内存的合并访问问题，内存合并访问几乎是现阶段所有GPU编程中非常重要而基础的一个技术细节等。计算密集型算法：要考虑如何减少指令数量等。

4、精度与性能

有些算法可以通过牺牲精度来提升性能。例如有些GPU的float16性能往往能翻倍，可以尝试用低精度的数据类型看看能否满足需要，cuda一样有这样的情况

5、使用fast-math编译选项

cuda内部有实现一些快速函数，可以尝试开启fast-math编译选项，或许有很明显的提升

6、使用release版本进行性能测试

cuda代码尽可能用release版本，对于在windows开发的同学应该深有感触，debug和relWithDebinfo的版本性能下降达到1个数量级

7、必备cuda相关开发工具

工欲善其事，必先利其器。cuda开发有比较完整的调试和优化工具能大大提高开发和调优效率，例如最新的有nsight compute，cuda debugger，nsight system。

作者：头像是狐狸吗
https://www.zhihu.com/question/437131193/answer/1652719667

写小东西的时候不难。因为算法简单，很容易发现哪些地方应该怎么优化。

(小东西指的是经典的卷积之类的教科书性能优化example。

但是一旦你写的东西真的复杂起来了。就要考虑很多东西。首先算法结构上，怎么拆依赖转并行。算法设计上就能让效率轻松拉开几个数量级。

然后是代码优化，你怎么布局线程。怎么分配/重用寄存器，怎么分配/对齐shared memory，怎么处理同步。怎么用好wave instructions。一个极限优化的程序和一个随便写写的程序会有极端情况下几个数量级的差距。

但是有时候这些东西真的不是普通人脑能分析清楚的。一直期待能有一个更高层的可以自动做这些优化的语言出现。之前很看好太极，然后有点失望。

反正现在我基本处于自暴自弃的状态，并不指望任何时候都能写出较优的实现。

作者：hello
https://www.zhihu.com/question/437131193/answer/1660230783

难在cuda是众核编程，有几千个核要调度。

拿当下最主流的张量计算来说，一个张量的输入是4维，输出也是4维，这4维的大小是不确定的，可能是3*3*512*1024，或者128*128*80*16，或者16*16*1024*256，等等。不同的维度（未经优化时）提供的并行能力是不一样的，你需要组织几千个线程，使他们在这些不同的维度上都尽可能的忙碌起来。

举例来说，输入128*128*3*256、输出1*1*3*256，和输入16*16*1024*256、输出16*16*1*256，代码上就是两套实现，kernel launch时也要仔细分配grid和block。

同时，一个函数的优化还要适配其他参数，比如为不同的卷积核优化，比如不同的layout，这就更加复杂，完整的优化一遍工作量很大很麻烦。

这还没算上不同的算法实现、不同的硬件架构。

作者：职业失败人士
https://www.zhihu.com/question/437131193/answer/1674334326

一个是并行程序本身的性质

对于不涉及同步的并行程序来说是一般是不难的，比如给数组每一个元素＋1。因为这种问题可以拆分成若干个互不干扰的等价于原问题的子问题，而每一个子问题可以当做一个串行程序来写。

而对于涉及同步的程序来说，比如排序，你需要考虑多个线程之间的同步和合作的问题，这样的程序写起来是没有上一种直观的。

另外一个问题是CUDA的抽象程度

CUDA为了给你细粒度控制程序内存分配与线程协作的方式给你抽象出了grid block warp thread的层级，同时给你了shared memory这种层级的内存，还有warp级别的指令。这相比只有thread的概念的情况下对程序的组织更加不直观。

最后就是性能问题

Nvidia的GPU架构众多，访存资源上从shared memory到寄存器文件差别很大，计算资源从tensor core到CUDA core的数目差别也很大。更不用说一些架构的特性比如bank conflict, memory coalescing等。你很难直接写出在大部分架构上高效的程序。

☆ END ☆