CUDA：爲什麼compute_20代碼在compute_35設備上失敗？

Question

對於泰坦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) 

的代碼編譯沒有任何錯誤，我的電腦上的巨人。但是，當我（我的電腦泰坦再次）運行它，它的一個thrust::copy呼叫，出現以下錯誤失敗後：

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

谷歌說，上述錯誤是由於GPU架構不匹配造成的。

最奇怪的部分是，上面的行（arch=compute_20,code=sm_21），代碼編譯和運行沒有錯誤在我的朋友的電腦與GTS 450！除GPU之外，她的Ubuntu 12.04，gcc和CUDA SDK 5.5版本與我的相同。

這是這個錯誤的真正原因？爲什麼泰坦不能運行compute_20的代碼？ CUDA GPU是不是應該與PTX或SASS代碼向後兼容？即使不是，爲什麼驅動程序JIT不能將compute_20 PTX編譯爲sm_35的SASS？

Answer 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，並直接非JIT執行上cc2.1或cc3.5設備。