title: SASS code and its corresponding asm code in kernel
body: <p>I have written the following one line cuda PTX instruction using <code>asm volatile</code></p>
  
  <pre><code>__global__ void add( uint32_t a, uint32_t b )
  {
    uint32_t c = 0;
    asm volatile("add.u32 %0, %1, %2;" : "=r"(c) : "r"(a), "r"(b) );
    printf("sink=%d\n", c);
  }
  </code></pre>
  
  <p>With this kernel call</p>
  
  <pre><code>add&lt;&lt;&lt; 1, 1 &gt;&gt;&gt;( 1,2 );
  </code></pre>
  
  <p>and this compilation command</p>
  
  <pre><code>nvcc -arch=sm_70 -Xptxas -O3,-v,-dlcm=ca -o test1 test1.cu
  </code></pre>
  
  <p>I see this sass code</p>
  
  <pre><code>    code for sm_70
              Function : _Z6addjj
      .headerflags    @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
      /*0000*/                   IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;   /* 0x00000a00ff017624 */
                                                                           /* 0x000fd000078e00ff */
      /*0010*/              @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ ;             /* 0x000000fffffff389 */
                                                                           /* 0x000fe200000e00ff */
      /*0020*/                   IMAD.MOV.U32 R0, RZ, RZ, c[0x0][0x160] ;  /* 0x00005800ff007624 */
                                                                           /* 0x000fe200078e00ff */
      /*0030*/                   IADD3 R1, R1, -0x8, RZ ;                  /* 0xfffffff801017810 */
                                                                           /* 0x000fe40007ffe0ff */
      /*0040*/                   MOV R4, 0x0 ;                             /* 0x0000000000047802 */
                                                                           /* 0x000fe40000000f00 */
      /*0050*/                   IADD3 R0, R0, c[0x0][0x164], RZ ;         /* 0x0000590000007a10 */
                                                                           /* 0x000fe40007ffe0ff */
      /*0060*/                   IADD3 R6, P0, R1, c[0x0][0x20], RZ ;      /* 0x0000080001067a10 */
                                                                           /* 0x000fe40007f1e0ff */
      /*0070*/                   MOV R5, 0x0 ;                             /* 0x0000000000057802 */
                                                                           /* 0x000fc60000000f00 */
      /*0080*/                   IMAD.X R7, RZ, RZ, c[0x0][0x24], P0 ;     /* 0x00000900ff077624 */
                                                                           /* 0x000fe200000e06ff */
      /*0090*/                   STL [R1], R0 ;                            /* 0x0000000001007387 */
                                                                           /* 0x0001ee0000100800 */
      /*00a0*/                   MOV R20, 0x0 ;                            /* 0x0000000000147802 */
                                                                           /* 0x000fe40000000f00 */
      /*00b0*/                   MOV R21, 0x0 ;                            /* 0x0000000000157802 */
                                                                           /* 0x000fd00000000f00 */
      /*00c0*/                   CALL.ABS.NOINC 0x0 ;                      /* 0x0000000000007943 */
                                                                           /* 0x001fea0003c00000 */
      /*00d0*/                   EXIT ;                                    /* 0x000000000000794d */
                                                                           /* 0x000fea0003800000 */
      /*00e0*/                   BRA 0xe0;                                 /* 0xfffffff000007947 */
                                                                           /* 0x000fc0000383ffff */
      /*00f0*/                   NOP;                                      /* 0x0000000000007918 */
                                                                           /* 0x000fc00000000000 */
              ......................
  </code></pre>
  
  <p>I know that compiler optimizes things, however, I expect to see 2 loads, 1 add and 1 store.</p>
  
  <p>What should I do in order to see that? Also, which sass instruction exactly corresponds to my <code>add.u32</code> instruction?</p>