CUDA程序优化分析和总结.docxVIP

CUDA 程序优化 确定任务中的串行和并行部分，选择合适的算法，将问题分为几个步骤看那些可以用并行来实现，确定要使用的算法 按照算法确定任务和数据的划分方式，将每个需要并行实现的步骤映射为一个 满足 CUDA 两层并行模式的内核函数，要尽量让每个 sm 上拥有至少六个活动的 warp 核至少两个活动的线程块。 编写一个能正常运行的程序作为优化的起点 优化显存的访问，避免显存的带宽成为瓶颈。在显存带宽优化前，其他的优化不会产生明显的结果。 显存优化的主要方法 将可以采用相同的 block 和 grid 的维度实现几个 kernel 的合并减少对显存的访问 尽量不要让私有变量分配到 local memory 为满足合并访问，采用 cudaMallocPitch()或者 cudaMalloc3D()分配显存。 为满足合并访问，将数据进行对齐（使用 align） 为满足合并访问，保证访问的首地址从 16 的整数倍开始，如果可能， 尽量让每个线程一次读的数据字长都为 32bit； 在数据只会被访问一次的，并且满足合并访问时，考虑使用 zerocopy 在某些情况下，考虑存储器控制器负责不均衡造成的风情冲突 使用拥有换成的常数存储器和纹理存储器 优化指令流 如果只需要少量的线程进行操作，一定要使用 类似‘if threaded N’ 的方式，避免多个线程同时运行占用更多的时间或者产生错误的结果 在不会出现不可接受的误差的情况下，采用 CUDA 算术指令集中的快速指令 使用#unroll 让编译器有效的展开循环 采用原子函数实现更复杂的算法，并保证结果的正确性 避免多余的同步 如果不产生 bank conflict 的算法不会造成算法的效率的的下降或者非合并访问，究应该避免 bank conflict 资源均衡 调整 shared memory 的使用量和 register 的使用量，保证更高的 SM 占用率，调整每个线程的数据处理量，shared memory，register 的使用量 使用括号明确变量的生存周期，使用 shared memory 存储变量 b) 与主机之间的通信的优化 尽量减少通信 使用 cudaMallocHost()分配主机的内存，获得更大的带宽 一次缓存较多的数据，在一并传输，可以获得较大的带宽 用流和异步处理来隐藏通信时间 使用 zerocopy 和 write-combined memory 提高可用带宽