建议和最佳做法:
整体性能优化策略:
性能优化围绕三个基本策略展开:
- 最大限度地平行执行
- 优化内存使用量以实现最大内存带宽
- 优化指令使用率以实现最大指令吞吐量
最大化并行执行从构建算法开始,尽可能多地暴露数据并行。 一旦算法的并行性暴露出来,它就需要尽可能有效地映射到硬件。 这是通过仔细选择每个内核启动的执行配置来完成的。 应用程序还应该通过流显式公开设备上的并发执行以及最大化主机和设备之间的并发执行,从而最大限度地提高并行执行的水平。
优化内存使用开始于最小化主机和设备之间的数据传输,因为这些传输比内部设备数据传输具有更低的带宽。 内核对全局内存的访问也应通过最大限度地使用设备上的共享内存来最小化。 有时候,最好的优化甚至可以首先通过简单地重新计算数据来避免任何数据传输。
根据每种存储器的访问模式,有效带宽可以变化一个数量级。 因此,优化内存使用的下一步是根据最佳内存访问模式来组织内存访问。 这种优化对于全局内存访问尤其重要,因为访问延迟需要花费数百个时钟周期。 共同的内存访问通常只有在存在高度银行冲突时才值得优化。
至于优化指令的使用,应该避免使用吞吐量低的算术指令。 这表明在不影响最终结果时速度的交易精度,例如使用内在函数而不是常规函数或单精度而不是双精度。 最后,由于设备的SIMT(单指令多线程)性质,必须特别注意控制流程指令。
NVCC编译器开关:
nvcc:
NVIDIA nvcc编译器驱动程序将.cu文件转换为C,用于主机系统和CUDA程序集或设备的二进制指令。 它支持许多命令行参数,其中以下对于优化和相关的最佳实践特别有用:
‣-maxrregcount = N指定内核在每个文件级别可以使用的最大寄存器数量。 见注册压力。 (另请参阅CUDA C编程指南的“执行配置”中讨论的__launch_bounds__限定符,以控制每个内核基础上使用的寄存器数量。)
‣--ptxas-options = -v或-Xptxas = -v列出每个内核的寄存器,共享和常量内存使用情况。
‣-ftz = true(非规格化数字被刷新为零)
‣-prec-div = false(不太精确的划分)
‣-prec-sqrt = false(不精确的平方根)
n -use_fast_math nvcc的编译器选项强制每个functionName()调用等效的__functionName()调用。 这使得代码运行速度更快,代价是精度和准确性降低。 请参阅数学库。