2

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

4

2 に答える 2

2

PTX は、コンパイルされたコードの中間表現にすぎません。GPU が実際に実行するものではありません。GPU が実行するコードを生成するアセンブリ ステップがさらにあります。これは、コンパイル時、またはドライバーで JIT コンパイルを使用して発生する可能性があります。その結果、あなたの指示は重要であり、あなたがそれらから推測したものはすべて無効です。

NVIDIAcuobjdumpは、Fermi カード用に生成されたアセンブラー出力を逆アセンブルし、GPU で実行される実際のマシン コードを表示できるツールを出荷しています。

于 2011-07-06T08:23:11.083 に答える
2

PTX は、GPU で実行されているものとまったく同じではないことに注意してください。PTX は単なる中間表現です。実際のコードは .cubin ファイルにあります。そのため、ptx ソース コードに基づいてそのような正確な計算を行っても意味がありません。

CUDA 4.0 で提供されるツールを使用cuobjdump --sassして、.cubin ファイルから GPU アセンブリ コードをより読みやすい形式に抽出できます。

于 2011-07-06T08:25:06.857 に答える