更改 CUDA 中的 arch 参数使我使用更多寄存器答案

【问题标题】：Changing the arch argument in CUDA makes me use more registers更改 CUDA 中的 arch 参数使我使用更多寄存器
【发布时间】：2013-02-25 23:40:41
【问题描述】：

我一直在我的 Tesla K20m 上编写内核，当我使用 -Xptas=-v 编译软件时，我得到以下结果：

ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z9searchKMPPciPhiPiS1_' for 'sm_10'
ptxas info    : Used 8 registers, 80 bytes smem, 8 bytes cmem[1]

如您所见，只使用了 8 个寄存器，但是，如果我提到参数 -arch=sm_35，我的内核执行时间会显着增加，并且使用的寄存器数量也会增加，我想知道为什么

nvcc mysoftware.cu -Xptxas=-v -arch=sm_35 
ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z9searchKMPPciPhiPiS1_' for 'sm_35'
ptxas info    : Function properties for _Z9searchKMPPciPhiPiS1_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 21 registers, 16 bytes smem, 368 bytes cmem[0]

由于在多本书中都提到使用正确的卡架构可以提高性能，我想知道为什么我的卡会急剧下降。

谢谢。

编辑：类似问答：Registers and shared memory depending on compiling compute capability?

【问题讨论】：

编译器似乎做出了不同的优化决策。你能提供你所看到的简短的可编译复制器吗？我对资源使用的变化没有你所报告的急剧放缓那么感兴趣。
所以基本上，当我使用 'sm_10' 时，代码需要 2.000235 毫秒才能执行，但是当使用 'sm_35" 时，代码需要 2.538336 到 3.xxxxxx 毫秒，这在这个内核的情况下令人印象深刻。
我现在实际上不能给出代码，它使用了一部分不是我的代码。
可能是编译器错误，但如果没有源代码，则无法确定。尝试指定__launch_bounds__（尝试不同的值），看看是否有什么不同。
我实际上使用了定义数量的线程和固定数量的块，所以尝试 launch_bounds 是否重要？

标签： cuda nvidia cpu-registers

【解决方案1】：

使用 sm_20 及更高版本进行编译可实现 IEEE 数学和 ABI 合规性。这两个选项会增加寄存器数量并降低性能。这两个选项可以禁用。

【讨论】：

这两个通常是主要原因。编译器也意识到较新的架构提供了更多的寄存器，因此可能会故意使用它们来提高性能（例如，使用临时变量来缓存和中间结果而不是重新计算，或者更积极地提前安排加载通常我不建议关闭任何一个IEEE 合规性或 ABI。尤其是使用 ABI 可以启用许多较新的 CUDA 语言功能。除非存在溢出或占用的特定问题，否则程序员不应沉迷于寄存器的使用。
其实我对寄存器并不是很着迷，我的问题是如果我使用选项 -arch=sm_35 进行编译，时间会增加很多。 @greg 你会建议禁用这些选项吗？还是根本没有？或者继续使用 -arch=sm_10 编译我的软件，因为时间非常重要。
在这里找到了类似的答案：stackoverflow.com/questions/6038221/…
尝试配合 -ftz=true -prec-div=false -prec-sqrt=false。这会将单精度浮点数学配置为与 sm_1x 编译所使用的设置大致相同。如果您在 64 位平台上，增加寄存器使用的另一个原因可能是对于 sm_20 及更高版本，设备代码中的 64 位指针只能在非常有限的情况下优化为 32 位操作，而这种优化可以应用于 sm_1x 设备上的所有设备指针，因为它们都不支持超过 4GB 的内存。