如何理解下面的asm？

Question

__global__ void TEST_prog(int *data_in1, int *data_in2, int *data_out) // employing IF functions

{

unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x;

data_out[tid] = data_in1[tid] + data_in2[tid];

}

code for sm_10
    Function : _Z9TEST_progPiS_S_
.headerflags    @"EF_CUDA_SM10 EF_CUDA_PTX_SM(EF_CUDA_SM10)"

    /*0000*/        MOV.U16 R0H, g [0x1].U16;          /* 0x0023c78010004205 */
    /*0008*/        I2I.U32.U16 R1, R0L;               /* 0x04000780a0000005 */
    /*0010*/        IMAD.U16 R0, g [0x6].U16, R0H, R1; /* 0x0020478060014c01 */
    /*0018*/        SHL R2, R0, 0x2;                   /* 0xc410078030020009 */
    /*0020*/        IADD32 R0, g [0x4], R2;            /* 0x2102e800         */
    /*0024*/        IADD32 R3, g [0x6], R2;            /* 0x2102ec0c         */
    /*0028*/        GLD.U32 R1, global14[R0];          /* 0x80c00780d00e0005 */
    /*0030*/        GLD.U32 R0, global14[R3];          /* 0x80c00780d00e0601 */
    /*0038*/        IADD32 R1, R1, R0;                 /* 0x20008204         */
    /*003c*/        IADD32 R0, g [0x8], R2;            /* 0x2102f000         */
    /*0040*/        GST.U32 global14[R0], R1;          /* 0xa0c00781d00e0005 */
    ...................................

g [0x1] 存储了什么？

g[xx]是共享内存，什么时候设置的值？

ROH是R0的高16位？

ROL之前没有赋值，但是第二个inst读取了它？

我猜g [0x4]，g [0x6]，g [0x4]是内核args，但是为什么args设置在共享内存中？

Answer 1

关键信息是

EF_CUDA_SM10 EF_CUDA_PTX_SM(EF_CUDA_SM10)

即此代码是适用于计算能力 1.0 设备的 SASS。在计算 1.x 设备中，内核参数存储在共享内存中。在计算 2.x 设备中，为此目的引入了专用的常量内存库。

所以你对代码的阅读是完全正确的。