title: CUDA compiler produce unoptimal assembler
body: <p>I have compiled flowing simple test kernel (CUDA5, sm2.0):</p>
  
  <pre><code>__device__ void TestKernel(int *pdata)
      {
        int a0,b0,c0;
  
        a0 = pdata[0];
        b0 = pdata[1];
  
        c0 = a0 + b0;
        pdata[2] = c0;
      }
  </code></pre>
  
  <p>and expect something like flowing assembler:</p>
  
  <pre><code>LD R3,[R0]
  LD R4,[R0+4]
  IADD R4,R4,R3
  ST [R0+8],R4
  </code></pre>
  
  <p>but from cuobjdump --dump-sass I see flowing much longer result:</p>
  
  <pre><code>/*0000*/     /*0x10001de428000000*/     MOV R0, R4;
  /*0008*/     /*0x00001de428000000*/     MOV R0, R0;
  /*0010*/     /*0x00001de428000000*/     MOV R0, R0;
  /*0018*/     /*0x00001de428000000*/     MOV R0, R0;
  /*0020*/     /*0x0000dc8580000000*/     LD R3, [R0];
  /*0028*/     /*0x0c00dde428000000*/     MOV R3, R3;
  /*0030*/     /*0x10011c034800c000*/     IADD R4, R0, 0x4;
  /*0038*/     /*0x10011de428000000*/     MOV R4, R4;
  /*0040*/     /*0x00411c8580000000*/     LD R4, [R4];
  /*0048*/     /*0x10011de428000000*/     MOV R4, R4;
  /*0050*/     /*0x1030dc0348000000*/     IADD R3, R3, R4;
  /*0058*/     /*0x20001c034800c000*/     IADD R0, R0, 0x8;
  /*0060*/     /*0x00001de428000000*/     MOV R0, R0;
  /*0068*/     /*0x0000dc8590000000*/     ST [R0], R3;
  /*0070*/     /*0x00001de790000000*/     RET;
  /*0078*/     /*0x00001de780000000*/     EXIT;
  /*0080*/     /*0x00001de780000000*/     EXIT;
  </code></pre>
  
  <p>Very strange to me MOVs instruction in addresses 8,10,18,28,38,60
  also the immediate offset in load/store instruction doesn't used.
  So instead expected 4 (actually 6 including RET,EXIT) instruction I get 15
  What is possible reason?</p>