在費米GPU（GTX 580）異形執行的指令併發出指令奇怪的結果

Question

我的內核有PTX的版本是這樣的：

.version 2.2 
.target sm_20, texmode_independent 

.entry histogram(
     .param .u32 .ptr .global .align 4 histogram_param_0, 
     .param .u32 .ptr .global .align 4 histogram_param_1 
) 
{ 
     .reg .f32  %f<2>; 
     .reg .s32  %r<12>; 

_histogram: 
     mov.u32   %r1, %tid.x; 
     mov.u32   %r2, %envreg3; 
     add.s32   %r3, %r1, %r2; 
     mov.u32   %r4, %ctaid.x; 
     mov.u32   %r5, %ntid.x; 
     mad.lo.s32  %r6, %r4, %r5, %r3; 
     shl.b32   %r7, %r6, 2; 
     ld.param.u32 %r8, [histogram_param_0]; 
     add.s32   %r9, %r8, %r7; 
     ld.param.u32 %r10, [histogram_param_1]; 
     ld.global.f32 %f1, [%r9]; 
     add.s32   %r11, %r10, %r7; 
     st.global.f32 [%r11], %f1; 
     ret; 
} 

我爲我數了一下，只有13我的內核指令（未包括ret指令）。當我將工作項數設置爲5120時，工作組大小爲64.因爲有16個SM，每個SM中有32個標量處理器，所以上述代碼將在SM中執行10次。正如我預料的那樣，執行指令的數量應該是10 * 13 = 130.但是在我進行配置後，結果是：發佈的指令= 130，執行的指令= 100。 1.爲什麼發出的指令數與執行的指令數不同？沒有分支，所以他們不應該是平等的？ 2.爲什麼執行指令的數量比預期的少？是否應該至少執行ptx版本中的所有指令？ 3.高速緩存未命中（L1和L2）對發佈的指令數量和執行的指令數量有任何影響嗎？ 謝謝

Answer 1

PTX只是編譯代碼的中間表示。這不是GPU實際執行的內容。還有一個組裝步驟，它發出GPU運行的代碼，這可以在編譯時發生，也可以在驅動程序中使用JIT編譯。因此，你的指令很重要，你從他們推斷出的任何內容都是無效的。

NVIDIA推出一款名爲cuobjdump工具，它可以分解爲費米顯卡產生的彙編輸出並顯示在GPU上

Answer 2

實際的機器代碼運行記住PTX不被在GPU上執行的到底是什麼。 PTX僅僅是一種中間表示。真正的代碼在.cubin文件中。這就是爲什麼使用基於ptx源代碼的這種令人滿意的計算沒有意義的原因。

您可以使用CUDA 4.0提供的cuobjdump --sass工具從.cubin文件中提取GPU彙編代碼，使其更易讀。