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

【问题标题】：How can I get the nvcc CUDA compiler to optimize more?如何让 nvcc CUDA 编译器进行更多优化？
【发布时间】：2017-09-28 03:20:58
【问题描述】：

当使用 C 或 C++ 编译器时，如果我们通过 -O3 开关，执行会变得更快。在 CUDA 中，有没有等价的东西？

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

【问题讨论】：

【解决方案1】：

警告：使用 nvcc -O3 filename.cu 编译只会将 -O3 选项传递给宿主代码。

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

nvcc -Xptxas -O3,-v filename.cu

将要求对 cuda 代码进行优化级别 3（这是默认设置），而 -v 要求进行详细编译，它报告了非常有用的信息，我们可以考虑进一步优化技术（稍后会详细介绍）。

另一个可用于 nvcc 编译器的速度优化标志是 -use_fast_math，它将以牺牲浮点精度为代价使用内部函数（请参阅 Options for Steering GPU code generation）。

无论如何，根据我的经验，这种自动编译器优化选项通常不会带来很大的提升。通过显式编码优化可以实现最佳性能，例如：

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

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

【讨论】：

-Xptxas -O2 只是将优化级别设置为 2 的一半，-Xcompiler -O2 是另一个。 ptxas 和 cicc 都在优化编译器，每个编译器都有自己的（尽管部分重叠）优化集。
ptxas 的最大优化级别是多少？ -O4、-O5、-O6 等没有错误。

【解决方案2】：

nvcc 支持许多类似于以 CPU 为目标的 C/C++ 编译器的选项。这记录在the nvcc documentation；您还可以运行nvcc --help 来获取这些选项的详细说明（也许nvcc --help | less 可以更轻松地滚动浏览它们）。

默认优化级别实际上是-O3（除非您指定-G，用于调试，这会禁用大多数优化）。您可以改为指定 -O0 或 -O1 等，但这只会禁用优化。

如果您只想优化将在 GPU 上运行的代码，而不是将在 CPU 上运行的代码，您可以传递不同的优化开关directly to the ptxas device code compiler。

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

【讨论】：