title: What's the efficient way to swap two register variables in CUDA?
body: <p>I'm starting to write some CUDA code, and I want to do the equivalent of <code>std::swap()</code> for two variables within a kernel; they're in the register file (no spillage, not in some buffer, etc.). Suppose I have the following device code:</p>
  
  <pre><code>__device__ foo(/* some args here */) {
  
      /* etc. */
  
      int x = /* value v1 */;
      int y = /* value v2 */;
  
      /* etc. */
  
      swap(x,y);
  
      /* etc. */
  }        
  </code></pre>
  
  <p>Now, I could just write</p>
  
  <pre><code>template &lt;typename T&gt; void swap ( T&amp; a, T&amp; b )
  {
    T c(a); a=b; b=c;
  }
  </code></pre>
  
  <p>but I wonder - isn't there some CUDA built-in for this functionality?</p>
  
  <p><strong>Notes:</strong></p>
  
  <ul>
  <li>Yes, I want this to run for all threads.</li>
  <li>Never mind about whether I have enough registers or not. Assume that I have them.</li>
  </ul>