在拆卸CUDA微顯然冗餘操作

我有以下的內核進行了全球存儲矩陣out全局存儲器矩陣in的一個簡單的任務：在拆卸CUDA微顯然冗餘操作

__global__ void simple_copy(float *outdata, const float *indata){ 

    int x = blockIdx.x * TILE_DIM + threadIdx.x; 
    int y = blockIdx.y * TILE_DIM + threadIdx.y; 

    int width = gridDim.x * TILE_DIM; 

    outdata[y*width + x] = indata[y*width + x]; 

}

我檢查由cuobjdump傾倒拆卸微：

Function : _Z11simple_copyPfPKf 
/*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
/*0008*/  /*0x80001de218000000*/  MOV32I R0, 0x20;   R0 = TILE_DIM 
/*0010*/  /*0x00001c8614000000*/  LDC R0, c [0x0] [R0];  R0 = c 
/*0018*/  /*0x90009de218000000*/  MOV32I R2, 0x24;   R2 = 36 
/*0020*/  /*0x00209c8614000000*/  LDC R2, c [0x0] [R2];  R2 = c 

int x = blockIdx.x * TILE_DIM + threadIdx.x; 
/*0028*/  /*0x9400dc042c000000*/  S2R R3, SR_CTAid_X;   R3 = BlockIdx.x 
/*0030*/  /*0x0c00dde428000000*/  MOV R3, R3;     R3 = R3 ??? 
/*0038*/  /*0x84011c042c000000*/  S2R R4, SR_Tid_X;   R3 = ThreadIdx.x 
/*0040*/  /*0x10011de428000000*/  MOV R4, R4;     R4 = R4 ??? 
/*0048*/  /*0x8030dca32008c000*/  IMAD R3, R3, 0x20, R4;  R3 = R3 * TILE_DIM + R4 (contains x) 

int y = blockIdx.y * TILE_DIM + threadIdx.y; 
/*0050*/  /*0x98011c042c000000*/  S2R R4, SR_CTAid_Y; 
/*0058*/  /*0x10011de428000000*/  MOV R4, R4; 
/*0060*/  /*0x88015c042c000000*/  S2R R5, SR_Tid_Y; 
/*0068*/  /*0x14015de428000000*/  MOV R5, R5; 
/*0070*/  /*0x80411ca3200ac000*/  IMAD R4, R4, 0x20, R5;  R4 ...     (contains y) 

int width = gridDim.x * TILE_DIM; 
/*0078*/  /*0x50015de428004000*/  MOV R5, c [0x0] [0x14];  R5 = c 
/*0080*/  /*0x80515ca35000c000*/  IMUL R5, R5, 0x20;   R5 = R5 * TILE_DIM  (contains width) 

y*width + x 
/*0088*/  /*0x14419ca320060000*/  IMAD R6, R4, R5, R3;  R6 = R4 * R5 + R3  (contains y*width+x) 

Loads indata[y*width + x] 
/*0090*/  /*0x08619c036000c000*/  SHL R6, R6, 0x2;    
/*0098*/  /*0x18209c0348000000*/  IADD R2, R2, R6;    
/*00a0*/  /*0x08009de428000000*/  MOV R2, R2;     R2 = R2 ??? 
/*00a8*/  /*0x00209c8580000000*/  LD R2, [R2];    Load from memory - R2 = 

Stores outdata[y*width + x] 
/*00b0*/  /*0x1440dca320060000*/  IMAD R3, R4, R5, R3;   
/*00b8*/  /*0x0830dc036000c000*/  SHL R3, R3, 0x2; 
/*00c0*/  /*0x0c001c0348000000*/  IADD R0, R0, R3;   R0 = R0 + R3 
/*00c8*/  /*0x00001de428000000*/  MOV R0, R0;     R0 = R0 ??? 
/*00d0*/  /*0x00009c8590000000*/  ST [R0], R2;    Store to memory 

/*00d8*/  /*0x40001de740000000*/  BRA 0xf0; 
/*00e0*/  /*0x00001de780000000*/  EXIT; 
/*00e8*/  /*0x00001de780000000*/  EXIT; 
/*00f0*/  /*0x00001de780000000*/  EXIT; 
/*00f8*/  /*0x00001de780000000*/  EXIT;

反彙編代碼的頂部或旁邊的評論是我自己的。

正如您所看到的，有一些明顯無用的操作，在註釋中標記爲???。實質上，它們是登記冊本身的動作。

我有那麼兩個以下問題：

如果他們是沒用的，我相信他們用地消耗計算時間。我可以通過刪除它們來優化反彙編的微碼嗎？
PTX文件可以在CUDA代碼中內聯。但是，PTX只是跨GPU可移植性所需的中間語言。我能以某種方式「內聯」一個優化的反彙編微碼嗎？

非常感謝您提前。

編輯：相同的代碼編譯在釋放模式FOR SM = 2.0

Function : _Z11simple_copyPfPKf 
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)" 
/*0000*/  MOV R1, c[0x1][0x100];   /* 0x2800440400005de4 */ 
/*0008*/  S2R R0, SR_CTAID.Y;    /* 0x2c00000098001c04 */ 
/*0010*/  S2R R2, SR_TID.Y;     /* 0x2c00000088009c04 */ 
/*0018*/  S2R R3, SR_CTAID.X;    /* 0x2c0000009400dc04 */ 
/*0020*/  S2R R4, SR_TID.X;     /* 0x2c00000084011c04 */ 
/*0028*/  MOV R5, c[0x0][0x14];    /* 0x2800400050015de4 */ 
/*0030*/  ISCADD R2, R0, R2, 0x5;   /* 0x4000000008009ca3 */ 
/*0038*/  ISCADD R3, R3, R4, 0x5;   /* 0x400000001030dca3 */ 
/*0040*/  SHL R0, R5, 0x5;     /* 0x6000c00014501c03 */ 
/*0048*/  IMAD R2, R0, R2, R3;    /* 0x2006000008009ca3 */ 
/*0050*/  ISCADD R0, R2, c[0x0][0x24], 0x2; /* 0x4000400090201c43 */ 
/*0058*/  ISCADD R2, R2, c[0x0][0x20], 0x2; /* 0x4000400080209c43 */ 
/*0060*/  LD R0, [R0];      /* 0x8000000000001c85 */ 
/*0068*/  ST [R2], R0;      /* 0x9000000000201c85 */ 
/*0070*/  EXIT ;       /* 0x8000000000001de7 */

編輯：相同的代碼編譯在釋放模式FOR SM = 2.1

Function : _Z11simple_copyPfPKf 
.headerflags @"EF_CUDA_SM20 EF_CUDA_PTX_SM(EF_CUDA_SM20)" 
/*0000*/  MOV R1, c[0x1][0x100];   /* 0x2800440400005de4 */ 
/*0008*/  NOP;        /* 0x4000000000001de4 */ 
/*0010*/  MOV R0, c[0x0][0x14];    /* 0x2800400050001de4 */ 
/*0018*/  S2R R2, SR_CTAID.Y;    /* 0x2c00000098009c04 */ 
/*0020*/  SHL R0, R0, 0x5;     /* 0x6000c00014001c03 */ 
/*0028*/  S2R R3, SR_TID.Y;     /* 0x2c0000008800dc04 */ 
/*0030*/  ISCADD R3, R2, R3, 0x5;   /* 0x400000000c20dca3 */ 
/*0038*/  S2R R4, SR_CTAID.X;    /* 0x2c00000094011c04 */ 
/*0040*/  S2R R5, SR_TID.X;     /* 0x2c00000084015c04 */ 
/*0048*/  ISCADD R2, R4, R5, 0x5;   /* 0x4000000014409ca3 */ 
/*0050*/  IMAD R2, R0, R3, R2;    /* 0x200400000c009ca3 */ 
/*0058*/  ISCADD R0, R2, c[0x0][0x24], 0x2; /* 0x4000400090201c43 */ 
/*0060*/  ISCADD R2, R2, c[0x0][0x20], 0x2; /* 0x4000400080209c43 */ 
/*0068*/  LD R0, [R0];      /* 0x8000000000001c85 */ 
/*0070*/  ST [R2], R0;      /* 0x9000000000201c85 */ 
/*0078*/  EXIT ;       /* 0x8000000000001de7 */

來源

2013-10-01 JackOLantern

這個CUDA的內置版本是什麼版本，以及編譯器選項用於構建代碼？ – njuffa

我懷疑上面的SASS不是來自發布版本，即代碼不是通過完全優化構建的。我使用CUDA 5.5 toolchain使用默認的編譯器設置編譯了上述各種體系結構的代碼，但沒有看到MOV指令。我將TILE_DIM定義爲32. – njuffa

編譯器可以將多行源代碼編譯成一行機器碼。如果在調試模式下發生這種情況，編譯器會爲沒有生成機器代碼的源代碼行插入額外的「no op」指令，以便爲這些行提供斷點位置。 –

答案這兩個問題都不是。

如果您嘗試刪除最終二進制有效內容中的指令。您將更改代碼段的長度並打破ELF和fatbinary文件。要解決這個問題，需要手工製作其格式不容易記錄的頭文件，這聽起來像是很多工作，只是爲了優化一些說明。

並不支持內聯本機彙編器，但我相信你已經知道了。

最後，我不能重現使用CUDA 5.0：

Fatbin elf code: 
================ 
arch = sm_20 
code version = [1,6] 
producer = cuda 
host = mac 
compile_size = 32bit 
identifier = pumpkinhead.cu 

    code for sm_20 
     Function : _Z11simple_copyPfPKf 
    /*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
    /*0008*/  /*0x98001c042c000000*/  S2R R0, SR_CTAid_Y; 
    /*0010*/  /*0x88009c042c000000*/  S2R R2, SR_Tid_Y; 
    /*0018*/  /*0x9400dc042c000000*/  S2R R3, SR_CTAid_X; 
    /*0020*/  /*0x84011c042c000000*/  S2R R4, SR_Tid_X; 
    /*0028*/  /*0x08001ca340000000*/  ISCADD R0, R0, R2, 0x5; 
    /*0030*/  /*0x10309ca340000000*/  ISCADD R2, R3, R4, 0x5; 
    /*0038*/  /*0x50001ca350004000*/  IMUL R0, R0, c [0x0] [0x14]; 
    /*0040*/  /*0x08009ca340000000*/  ISCADD R2, R0, R2, 0x5; 
    /*0048*/  /*0x90201c4340004000*/  ISCADD R0, R2, c [0x0] [0x24], 0x2; 
    /*0050*/  /*0x80209c4340004000*/  ISCADD R2, R2, c [0x0] [0x20], 0x2; 
    /*0058*/  /*0x00001c8580000000*/  LD R0, [R0]; 
    /*0060*/  /*0x00201c8590000000*/  ST [R2], R0; 
    /*0068*/  /*0x00001de780000000*/  EXIT; 
     .....................................

你確定你已經證明與發行設置編譯的代碼？

來源

2013-10-02 04:53:15 talonmies

非常感謝您的回答。事實上，代碼是在調試模式下編譯的。我在發佈模式下編譯時發佈了微代碼，代碼看起來與昨天發佈的代碼非常不同，與您的代碼非常相似。我不能完全複製你的，但也許這是由於我們正在使用的稍微不同的優化選項。 – JackOLantern

在拆卸CUDA微顯然冗餘操作

回答

相關問題