CUDA：为什么 compute_20 代码在 compute_35 设备上失败？答案

【问题标题】：CUDA: Why does compute_20 code fail on compute_35 device?CUDA：为什么 compute_20 代码在 compute_35 设备上失败？
【发布时间】：2016-02-13 16:12:59
【问题描述】：

对于带有 Titan GPU (compute_35,sm_35) 的计算机，我使用 CMakeLists.txt 中的这一行编译了一些代码：

set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_35,code=sm_35)

代码可以编译并且运行良好。

我想检查一下这段代码会给使用 GTS 450 (compute_20,sm_21) 的朋友带来哪些编译问题。所以，我将上面的行改为：

set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_20,code=sm_21)

代码在我的计算机上使用 Titan 编译时没有任何错误。但是当我运行它（再次在我的 Titan 计算机上）时，它在 thrust::copy 调用后失败，并出现以下错误：

$ ./foobar
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  invalid device function 
"foobar" terminated by signal SIGABRT (Abort)

Google 称上述错误是由于 GPU 架构不匹配造成的。

最奇怪的是，使用上面的行 (arch=compute_20,code=sm_21)，代码在我朋友的 GTS 450 计算机上编译并运行没有错误！除了GPU，她的Ubuntu 12.04、gcc和CUDA SDK 5.5版本和我一样。

这是导致此错误的真正原因吗？为什么 Titan 不能运行 compute_20 代码？ CUDA GPU 不应该向后兼容 PTX 或 SASS 代码吗？即使不是，为什么驱动JIT不能将compute_20 PTX编译成sm_35的SASS？

【问题讨论】：

标签： cuda cmake

【解决方案1】：

如果您指定：

-gencode arch=compute_20,code=compute_20

您的代码应该在任一 GPU 上运行（通过 JIT）。

根据nvcc manual，当你为code开关指定虚拟架构时，直接启用JIT。您可以在一个命令中进行多个规范：

-arch=compute_20 -code=compute20,sm_21,sm_35

（注意这是代替指定-gencode ...）

这将允许来自 sm_20 PTX 的 JIT，以及直接在 cc2.1 或 cc3.5 设备上执行非 JIT。

【讨论】：

我没有意识到compute_xx 可以传递给-code！谢谢:)