CUDA SM10架构下TEST_prog内核汇编代码相关技术问询

原始CUDA内核代码 #

__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];
}

sm_10架构对应的PTX汇编代码 #

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 */
    ...................................

问题与解答 #

• 问题1：g [0x1]中存储的是什么内容？
  g[0x1]存储的是当前线程所属线程块的索引blockIdx.x（16位数值），后续会用它参与计算线程全局索引tid。

• 问题2：g[xx]为共享内存，其值何时被设置？
  这些共享内存中的值是在内核启动、线程块开始执行前，由CUDA硬件自动初始化设置的。包括线程/块索引、块维度等执行环境参数，以及内核参数，都会在硬件调度线程块时被加载到共享内存的固定位置。

• 问题3：R0H是否为寄存器R0的高16位？
  是的。PTX汇编语法中，后缀H代表32位寄存器的高16位，L代表低16位，所以R0H就是寄存器R0的高16位部分。

• 问题4：第二条指令读取R0L，但此前未对其赋值，这是为何？
  在SM10架构中，线程启动时硬件会自动把当前线程的threadIdx.x值加载到R0寄存器的低16位（R0L），属于硬件层面的默认行为，不需要显式指令赋值，因此可以直接读取使用。

• 问题5：推测g [0x4]、g [0x6]、g [0x8]为内核参数，为何参数会被存放在共享内存中？
  早期SM10架构没有专门的参数寄存器传递内核参数，CUDA runtime会把内核参数复制到线程块的共享内存中。共享内存访问延迟远低于全局内存，线程块内所有线程都能快速读取这些参数，避免了每个线程单独从全局内存读取参数的性能开销。

内容的提问来源于stack exchange，提问作者user24077790