title: Bug in PTX ISA (carry propagation)?
body: <p>Is there a bug in Cuda? I have run the following code on my GTX580 and r1 is zero at the end. I expect that it is one due to carry propagation? I have tested the code with Cuda Toolkit 4.2.9 and 5.5 and use "nvcc -arch=sm_20 bug.cu -o bug  &amp;&amp; ./bug" to compile and run it.</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;cuda.h&gt;
  
  __global__ void bug()
  {
    unsigned int r1 = 0;
    unsigned int r2 = 0;
  
    asm( "\n\t"
         "sub.cc.u32 %0, 0, 1;\n\t"
         "addc.cc.u32 %1, 0, 0;\n\t"
       : "=r"(r1), "=r"(r2) );
  
    printf("r1 &gt;&gt; %04X\n", r1);
    printf("r2 &gt;&gt; %04X\n", r2);
  
  }
  
  int main(void)
  {
    float *a_d;
    cudaMalloc((void **) &amp;a_d, 1);
  
    bug &lt;&lt;&lt; 1,1 &gt;&gt;&gt; ();
  
    cudaFree(a_d);
  }
  
  Output
  r1 &gt;&gt; FFFFFFFF
  r2 &gt;&gt; 0000
  </code></pre>