1
我編流動簡單的測試內核(CUDA5,SM2.0):CUDA編譯器產生unoptimal彙編
__device__ void TestKernel(int *pdata)
{
int a0,b0,c0;
a0 = pdata[0];
b0 = pdata[1];
c0 = a0 + b0;
pdata[2] = c0;
}
,並希望像流淌彙編:
LD R3,[R0]
LD R4,[R0+4]
IADD R4,R4,R3
ST [R0+8],R4
但cuobjdump --dump - 我看到流動更長的結果:
/*0000*/ /*0x10001de428000000*/ MOV R0, R4;
/*0008*/ /*0x00001de428000000*/ MOV R0, R0;
/*0010*/ /*0x00001de428000000*/ MOV R0, R0;
/*0018*/ /*0x00001de428000000*/ MOV R0, R0;
/*0020*/ /*0x0000dc8580000000*/ LD R3, [R0];
/*0028*/ /*0x0c00dde428000000*/ MOV R3, R3;
/*0030*/ /*0x10011c034800c000*/ IADD R4, R0, 0x4;
/*0038*/ /*0x10011de428000000*/ MOV R4, R4;
/*0040*/ /*0x00411c8580000000*/ LD R4, [R4];
/*0048*/ /*0x10011de428000000*/ MOV R4, R4;
/*0050*/ /*0x1030dc0348000000*/ IADD R3, R3, R4;
/*0058*/ /*0x20001c034800c000*/ IADD R0, R0, 0x8;
/*0060*/ /*0x00001de428000000*/ MOV R0, R0;
/*0068*/ /*0x0000dc8590000000*/ ST [R0], R3;
/*0070*/ /*0x00001de790000000*/ RET;
/*0078*/ /*0x00001de780000000*/ EXIT;
/*0080*/ /*0x00001de780000000*/ EXIT;
對我很陌生MOVs指令在地址8,10,18,28 ,38,60 也不使用加載/存儲指令中的直接偏移量。 所以反而期望4(實際上6包括RET,EXIT)指令我得到15 什麼是可能的原因?
您顯示的代碼是設備函數,而不是內核,或者是一個錯誤?此外,你確定這不是一個帶有調試設置的版本(當然看起來像)。請添加您使用的確切編譯命令。 – talonmies
你用'-G'編譯?我只會得到類似你用'-G'編譯時發佈的代碼。如果沒有,我會得到一個包含RET的5個指令的函數。 –
-G應該是調試信息,而不是禁用優化,並且它對分析有用。我將重新檢查此問題,但 – BaraBashkaD