如何让nvcc CUDA编译器进行更多优化？

Question

使用C或C++编译器时,如果我们通过-O3交换机,执行会变得更快.在CUDA中,有没有相同的东西？

我正在使用该命令编译我的代码nvcc filename.cu.之后我执行./a.out.

Answer 1

警告:编译nvcc -O3 filename.cu将仅将-O3选项传递给主机代码.

为了优化CUDA内核代码,必须将优化标志传递给PTX编译器,例如:

nvcc -Xptxas -O3,-v filename.cu

将要求优化级别3到cuda代码(这是默认值),同时-v请求详细编译,它报告了我们可以考虑用于进一步优化技术的非常有用的信息(稍后将详细介绍).

可用于nvcc编译器的另一个速度优化标志-use_fast_math是以浮点精度为代价使用内在函数(请参阅引导GPU代码生成的选项).

无论如何,根据我的经验,这种自动编译器优化选项通常无法实现很大的提升.通过显式编码优化可以实现最佳性能,例如:

1. 指令级并行(ILP):让每个CUDA线程在多个元素上执行其任务 - 这种方法将保持管道加载并最大化吞吐量.例如,假设您要处理NxN磁贴的元素,您可以使用2级TLP启动NxM线程块(其中M = N/2)并让threadIdx.y循环遍历2个不同的元素行.
2. 寄存器溢出控制:控制每个内核使用的寄存器数量并试用该-maxrrregcount=N选项.内核需要的寄存器越少,就越有资格同时运行(直到寄存器溢出将接管).
3. 循环展开:尝试#pragma unroll N在CUDA内核中的任何独立循环(如果有)之前添加.N可以是2,3,4.当您在注册压力和实现展开水平之间达到良好平衡时,可以获得最佳结果.这种方法毕竟属于ILP技术.
4. 数据打包:有时您可以将不同的相关缓冲区数据连接float A[N],B[N]到一个float2 AB[N]数据缓冲区中.这将转化为负载/存储单元和总线使用效率的更少操作.

当然,始终始终检查您的代码,以便合并内存访问全局内存并避免共享内存中的库冲突.使用nVIDIA Visual Profiler可以更深入地了解此类问题.

Answer 2

nvcc支持许多与CPU目标C/C++编译器类似的选项.这在nvcc文档中有记录 ; 你也可以运行nvcc --help以获得这些选项的详细描述(也许nvcc --help | less能够更容易地滚动它们).

实际上,默认优化级别-O3(除非您指定-G,否则用于调试,禁用大多数优化).您也可以指定-O0或-O1等,但只会禁用优化.

如果您只想优化将在GPU上运行的代码,而不是将在CPU上运行的代码,则可以将不同的优化开关直接传递给ptxas设备代码编译器.

此外,如果您编写nvcc -o foo filename.cu生成的可执行文件将被命名foo而不是a.out,以防您需要可执行文件的有意义的名称.这也与C/C++编译器相同.