CODEWITHSUNDEEP
cudanvidia
Ginu Jacob
Ginu Jacob

Reputation: 1778

Thread ID rotation in NVIDIA GPU assembly code (SASS)

While having a closer look at the SASS output for NVIDIA sm_20 architechture, it was observed that the ThreadID is loaded from a special register and rotation is performed in 2 steps.

  1. Load MSB register with the value obtained by rotating ThreadID left by 2
  2. Load LSB register with the value obtained by rotating ThreadID right by 30

Together this does a 32 bit rotation and the register pair will look like: enter image description here

Further, this register pair is used in the SASS wherever threadID is used in the code. The SASS code is

    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

What would be the purpose of doing it this way instead of using the ThreadID loaded from the special register?

Upvotes: 1

Views: 287

Answers (1)

Greg Smith
Greg Smith

Reputation: 11529

code for sm_20 Function : _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

Upvotes: 1

Related Questions