CUDA 程序优化

一、优化显存访问

1.      将可以采用相同的block和grid维度实现的几个kernel合并为1个，减少对显存的访问。

2.      除非非常必要，应该尽力避免将线程私有变量分配到local memory。

3.      为满足合并访问，采用cudaMallocPith()或cudaMalloc3D()分配显存。

4.      为满足合并访问，对数据类型进行对齐（使用__align）。

5.      为满足合并访问，保证访问的首地址从16的整数倍开始，如果可能，尽量让每个进程一次读的数据字长为32bit。

6.      在数据只会被访问一次，并且满足合并访问的情况下，可以考虑使用zerocopy。

7.      使用拥有缓存的常量存储器和纹理存储器提高某些应用的实际带宽。

 

二、计算量

如果计算量太小，那么使用CUDA并不划算。

 

三、开发流程

1.      确定任务中的串行部分和并行部分，选择合适的算法。

首先，需要将问题分为几个步骤，然后确定哪些步骤可以用并行算法实现，并确定要使用的算法。

2.      按照算法确定数据和任务的划分方式，将每个需要并行实现的步骤映射为一个满足CUDA并行模型的内核函数。

在这里就要尽量让每个SM上拥有至少6个活动的warp(6*32 = 192个线程)和至少2个活动线程块。

四、优化指令流

1.      如果只需要少量线程进行操作，一定记得要使用类似“if threadID < N”的方式，避免多个线程同时运行占用更长时间或产生错误结果。

2.      避免多余的同步

3.      如果不产生bank conflict的算法不会产生造成算法效率的下降或者非法合并访问，就应该避免bank conflict

4.      采用原子操作实现更加复杂的算法，并保证结果的正确性。

五、资源均衡

调整sharedmemory和register的使用量。为了使程序能获得更高的SM占用率，应该调整每个线程处理的数据数量、shared memory 和register的使用量。

当线程处理的子任务间有一些完全相同的部分时，应该只使用少量线程来完成公共部分的计算，再将公用数据通过shared memory 广播给所有线程。

通过调整块block的大小，修改算法和指令，以及动态分配shared memory，都可以提高shared memory的利用率。

六、与主机通信优化

1.      一次缓存较多的数据，然后一并传输，可以获得较高的实际带宽

2.      需要将结果显示在屏幕时，直接使用与图形学API互操作功能完成，避免将数据返回。

3.      使用流或异步处理隐藏与主机的通信时间

4.      使用cudaMallocPith()或cudaMalloc3D()分配存储空间，使用cudaMemcpy2D()拷贝数据。