如何让 nvcc CUDA 编译器进行更多优化？
How can I get the nvcc CUDA compiler to optimize more?
当使用 C 或 C++ 编译器时，如果我们传递 -O3 开关，执行速度会更快。在 CUDA 中，是否有等效的东西？
我正在使用命令 nvcc filename.cu 编译我的代码。之后我执行 ./a.out.
nvcc 支持许多类似于 CPU-targeting 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++ 编译器相同。
警告： 使用 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 个不同的线程上循环元素线。
- register spilling control：控制每个内核使用的寄存器数量，并尝试使用
-maxrrregcount=N 选项。内核需要的寄存器越少，有资格同时 运行 的块就越多（直到寄存器溢出将接管）。
- 循环展开： 尝试在您的 CUDA 内核中的任何独立循环（如果有的话）之前添加
#pragma unroll N。 N可以是2、3、4。当您在套准压力和实现的展开水平之间达到良好平衡时，就会达到最佳效果。毕竟，这种方法属于ILP技术。
- 数据打包：有时您可以将不同的相关缓冲区数据（例如
float A[N],B[N]）合并到一个float2 AB[N]数据缓冲区中。这将转化为 load/store 单位的更少操作和总线使用效率。
当然，总是，总是，总是检查您的代码以合并对全局内存的内存访问并避免共享内存中的存储区冲突。使用 nVIDIA Visual Profiler 更深入地了解此类问题。
当使用 C 或 C++ 编译器时，如果我们传递 -O3 开关，执行速度会更快。在 CUDA 中，是否有等效的东西？
我正在使用命令 nvcc filename.cu 编译我的代码。之后我执行 ./a.out.
nvcc 支持许多类似于 CPU-targeting 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++ 编译器相同。
警告： 使用 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 个不同的线程上循环元素线。
- register spilling control：控制每个内核使用的寄存器数量，并尝试使用
-maxrrregcount=N选项。内核需要的寄存器越少，有资格同时 运行 的块就越多（直到寄存器溢出将接管）。 - 循环展开： 尝试在您的 CUDA 内核中的任何独立循环（如果有的话）之前添加
#pragma unroll N。 N可以是2、3、4。当您在套准压力和实现的展开水平之间达到良好平衡时，就会达到最佳效果。毕竟，这种方法属于ILP技术。 - 数据打包：有时您可以将不同的相关缓冲区数据（例如
float A[N],B[N]）合并到一个float2 AB[N]数据缓冲区中。这将转化为 load/store 单位的更少操作和总线使用效率。
当然，总是，总是，总是检查您的代码以合并对全局内存的内存访问并避免共享内存中的存储区冲突。使用 nVIDIA Visual Profiler 更深入地了解此类问题。