NVIDIA GPU 汇编代码中的线程 ID 旋转 (SASS)

Question

在仔细查看 NVIDIA sm_20 架构的 SASS 输出时，发现 ThreadID 是从一个特殊的寄存器加载的，并且分两步执行旋转。

用旋转获得的值加载MSB寄存器 ThreadID 剩下 2
用获得的值加载LSB寄存器通过将 ThreadID 向右旋转 30

一起进行 32 位旋转，寄存器对将如下所示：

此外，此寄存器对用于 SASS 代码中使用 threadID 的任何地方。 SASS代码是

    code for sm_20
            Function : _Z3addPiS_S_
    .headerflags    @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)"
    /*0000*/         MOV R1, c[0x1][0x100];         /* 0x2800440400005de4 */
    /*0008*/         S2R R2, SR_TID.X;              /* 0x2c00000084009c04 */---Getting thread ID
    /*0010*/         IMAD.U32.U32 RZ, R1, RZ, RZ;   /* 0x207e0000fc1fdc03 */
    /*0018*/         SHL.W R3, R2, 0x2;             /* 0x6000c0000820de03 */---Rotating Step 1
    /*0020*/         SHR.U32 R4, R2, 0x1e;          /* 0x5800c00078211c03 */---Rotating Step 2
    /*0028*/         IADD R6.CC, R3, c[0x0][0x20];  /* 0x4801400080319c03 */
    /*0030*/         IADD.X R7, R4, c[0x0][0x24];   /* 0x480040009041dc43 */
    /*0038*/         LD.E R0, [R6];                 /* 0x8400000000601c85 */
    /*0040*/         STS [R3], R0;                  /* 0xc900000000301c85 */---Shared mem access with ThreadID

这样做的目的是什么，而不是使用从专用寄存器加载的 ThreadID？

Answer 1

sm_20代码函数：_Z3addPiS_S_

Demangled = add(int*, int*, int*)

MOV R1, c[0x1][0x100];      
S2R R2, SR_TID.X;               // read threadIdx.x
IMAD.U32.U32 RZ, R1, RZ, RZ;
SHL.W R3, R2, 0x2;              // r3 = threadIdx.x * 4 (int* pointer math)
SHR.U32 R4, R2, 0x1e;           // r4 = threadIdx.x[31:30] to make 64-bit offset in R3/R4
IADD R6.CC, R3, c[0x0][0x20];   // add a constant (parameter 0 - lower 32-bits)
IADD.X R7, R4, c[0x0][0x24];    // add a constant (parameter 0 - upper 32-bits)
LD.E R0, [R6];                  // load the 32-bit value from address R6/R7 into R0
STS [R3], R0;                   // store the 32-bit value in R0 into shared offset threadIdx.x * 4

NVIDIA GPU 汇编代码中的线程 ID 旋转 (SASS)

Thread ID rotation in NVIDIA GPU assembly code (SASS)

cuda

nvidia