プロファイル実行される命令の奇妙な結果とフェルミGPU（GTX 580）の指示を出し

Question

私のカーネルは、このようなPTXバージョンがあります：私は数えてI

.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があり、それぞれに32個のスカラープロセッサがあるため、上記のコードはSMで10回実行されます。私が期待したように、実行される命令の数は10 * 13 = 130でなければなりません。しかし、私がプロファイリングした後、結果は発行された命令= 130、実行された命令= 100です。 1.発行された命令の数と実行された命令の数が異なるのはなぜですか？枝がないので、彼らは平等ではないと思われますか？ 2.実行された命令の数が予想よりも少ないのはなぜですか？ ptxバージョンのすべての命令を少なくとも実行する必要がありますか？ 3.キャッシュミス（L1およびL2）は、発行された命令の数および実行された命令の数に影響しますか？ ありがとう

Answer 1

PTXはコンパイルされたコードの中間表現です。 GPUが実際に実行するものではありません。 GPUが実行するコードを発行するアセンブリステップがさらにあります。これは、コンパイル時に、またはドライバでJITコンパイルを使用して実行できます。結果として、あなたの命令数とそれらから推測されるものは無効です。

は、NVIDIAは、PTXは、GPU上で実行されているまさにではないことに注意してくださいフェルミカード用に生成アセンブラ出力を分解し、GPU

Answer 2

上、実際のマシンコードの実行を表示することができますcuobjdumpと呼ばれるツールを出荷します。 PTXは単なる中間表現です。実際のコードは.cubinファイルにあります。 ptxソースコードに基づいてそのような正確な計算を行うのは、何の意味もありません。

CUDA 4.0付属のcuobjdump --sassツールを使用すると、.cubinファイルからGPUアセンブリコードをより読みやすいものに抽出できます。