NVIDIA GPU彙編代碼（SASS）中的線程ID旋轉

雖然仔細查看了NVIDIA sm_20架構的SASS輸出，但觀察到ThreadID是從特殊寄存器加載的，旋轉分兩步執行。通過旋轉獲得的值NVIDIA GPU彙編代碼（SASS）中的線程ID旋轉

負載MSB寄存器線程ID通過2
負載左LSB右30

一起但這32位旋轉通過旋轉線程ID獲得寄存器的值並且寄存器對將如下所示：

此外，該寄存器對用於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

什麼會做這種方式，而不是使用從特殊寄存器加載的線程ID的目的是什麼？

來源

2016-06-07 Ginu Jacob

什麼是整個內核函數在做什麼？也許你可以從更高層面得到線索。人們通常會在給出threadIdx.x的情況下計算出某些內容，然後將其用作訪問共享/全局內存的偏移量。 – kangshiyin

代碼sm_20 功能：_Z3addPiS_S_

Demangled =添加（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

來源

2016-06-08 05:45:54

我可以理解乘以4後的邏輯。這是因爲int的大小是4.但是當我們將R3-R4視爲寄存器對時，當2 MSB位放置在R4中時它有何幫助。 –

我將不得不看到源代碼。我懷疑編譯器將threadIdx.x（uint32_t）提升爲uint64_t，這將要求threadIdx.x的2個MSB向下移動到高32位的2個LSB。編譯器不知道threadIdx.x的範圍有限，因爲內置的數據類型是uint3。 –

NVIDIA GPU彙編代碼（SASS）中的線程ID旋轉

回答

相關問題