CUDA——性能优化（总结）

CUDA性能优化策略（总结）

1）每个块上的线程数

给定每网格的线程总数，设计每块的线程数或网格的块数时应该最大化可用计算资源的利用率。

a.块的数目
块的数目至少要大于等于你的多处理器（SM）的个数，这样才能充分调动所有多处理器。
每个多处理器上的块应存在两个或者以上的活动块。保证线程同步（块同步）时，有多的活动块可用调用。此时，每个块上的共享内存至多为每个多处理器上共享内存的一半。

b.线程的数目
有了足够大数目的块，每块线程的数目应为warp大小的倍数，为64的倍数是最好的选择（在一个warp堵塞时，可用有多个活动warp去调用）。但是，块上的线程数目不是越多越好，还得考虑寄存器的数目。若分配的线程数目越多，则可供每个线程用的寄存器越少，所以，我们要避免使用的寄存器数目超过上限，否则内核将会启动失败。

在CUDA项目编译时，可以添加--ptxas-options=-v,可以用于查看内核已经使用了多少个寄存器（以及本地、共享和常量内存的使用情况）。也可以用.ptx文件去查看。

2）宿主与设备端的数据传输

宿主内存与设备内存之间的带宽比较小，所以要尽量减少宿主与设备端之间的数据传输。
I) 有些中间数据可以在设备内存创建就在设备内存创建，由设备操作，销毁。

II) 将多次小传输合并成一次大传输。

III) 使用页锁定内存时，可以获得较大的宿主内存和设备内存带宽，再使用流管理，将宿主到设备之间的数据并发传输。

3）在适当的时候选择纹理内存

纹理内存有高速缓存，拥有更大的带宽。而且它不受访问模式的限制，如果是用全局内存的话，我们得合并去访问全局内存才能获得较大的带宽。
当进行简单的图像处理时（如resize图片，滤波等操作时）如果纹理绑定的是CUDA数组，可以拥有更高的性能。但缺点是无法对纹理内存进行写入操作。（详情可以查看此博客https://blog.csdn.net/weixin_44444450/article/details/104249441

整体优化策略（来自CUDA手册+个人理解）

性能优化围绕三个基本策略：

1.最大化并行执行；

最大化并行执行可以分为 线程级并行和任务级并行
线程级并行： 需要考虑算法的并行，如何去设计算法，使其更好的去达到一个并行的效果，尽可能有效地将其映射到硬件。通过仔细地选择内核的执行配置来完成此操作。
任务级并行： 指的是通过流的方式在设备上显示地并发执行，以获取较高级别的并行，如：宿主与设备之间的并发执行。

2.优化内存使用以获得最大内存带宽；

最大化内存应尽量不进行低带宽的数据传输，如：宿主内存到设备内存；设备的全局内存访问；
1）我们可以选用高带宽的内存，如：共享内存，常量内存，纹理内存。
2）提高内存的有效带宽，这个取决于访问模式，我们可以将同一个指令的warp合并起来去访问全局内存，以此提高全局内存的有效带宽。全局内存合并访问规则可以查看此博客https://blog.csdn.net/weixin_44444450/article/details/104401965

共享内存的优化的关键一点就是要避免存储体冲突（bank conflict）

3.优化指令使用以获得最大指令吞吐量；

应尽量使用高吞吐量的指令，可以将一些除运算和求余运算换成对应的位运算（当数量大的时候可以这么用，数量小的话可以不用，会降低代码的可阅读性）。
可以将一些代码进行循环展开，这可以有效地避免warp的分支冲突。
将一个warp的循环展开，可以不用添加__syncthreads()，因为warp里面的线程会自动同步。