title: CUDA: Why does compute_20 code fail on compute_35 device?
body: <p>For a computer with Titan GPU (<code>compute_35,sm_35</code>), I compiled some code using this line in <code>CMakeLists.txt</code>:</p>
  
  <pre><code>set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_35,code=sm_35)
  </code></pre>
  
  <p>The code compiles and also runs fine.</p>
  
  <p>I wanted to check what compilation problems this code would cause for a friend who uses a GTS 450 (<code>compute_20,sm_21</code>). So, I changed the above line to:</p>
  
  <pre><code>set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_20,code=sm_21)
  </code></pre>
  
  <p>The code compiles without any errors on my computer with Titan. But when I run it (again on my Titan computer), its fails after a <code>thrust::copy</code> call with the following error:</p>
  
  <pre><code>$ ./foobar
  terminate called after throwing an instance of 'thrust::system::system_error'
    what():  invalid device function 
  "foobar" terminated by signal SIGABRT (Abort)
  </code></pre>
  
  <p>Google says the above error is caused due to GPU architecture mismatch.</p>
  
  <p>The strangest part is that with the above line (<code>arch=compute_20,code=sm_21</code>), the code compiles and runs without error on my friend's computer with GTS 450! Except for the GPU, her Ubuntu 12.04, gcc and CUDA SDK 5.5 versions are the same as mine.</p>
  
  <p>Is this the real cause of this error? Why cannot Titan run <code>compute_20</code> code? Isn't a CUDA GPU supposed to be backwards compatible with PTX or SASS code? Even if it isn't, why cannot the driver JIT compile the <code>compute_20</code> PTX to the SASS of <code>sm_35</code>?</p>