1

GeForce GTX 550 Ti で 1 つの倍精度 FLOP を実行するのにかかるナノ秒数を取得することに興味があります。

これを行うために、私は次のアプローチに従っています。カードの単精度のピーク パフォーマンスは 691.2 GFLOPS であることがわかりました。これは、倍精度のピーク パフォーマンスがその 1/8、つまり 86.4 GFLOPS になることを意味します。次に、コアあたりの FLOPS を取得するために、86.4 GFLOPS をコア数 192 で割ると、コアあたり 0.45 GFLOPS が得られます。0.45 GFLOPS は、コアあたりナノ秒あたり 0.45 FLOPS を意味します。正しいアプローチに従っている場合、これらの GFLOPS 数値を取得するために実行されるコアあたりのスレッド数と、この情報はどこで確認できますか?

さらに、以下に示す小さなテストは、1 つのスレッドだけで 236000232 サイクルで実行されます。ループの 1 回の反復を実行するのにかかる時間 (ナノ秒単位) を見つけるために、236000232/10^6 = 236 サイクルを実行します。カードのシェーダー クロックは 1800Mhz です。これは、ループの 1 回の反復を実行するのに 236/1.8 = 131 ナノ秒かかることを意味します。この数値は上記の数値よりもはるかに大きい (コアあたり 0.45 ナノ秒)。数値が大きく異なるため、ここで何かが欠けていると確信しています。その背後にある数学を理解するのを手伝ってください。

    __global__ void bench_single(float *data)
{
    int i;
    double x = 1.;
    clock_t start, end;
    start = clock();
    for(i=0; i<1000000; i++)
    {
        x = x * 2.388415813 + 1.253314137;
    }
    end = clock();
    printf("End and start %d - %d\n", end, start);
    printf("Finished in %d cycles\n", end-start);
}

ありがとうございました、

4

2 に答える 2

4

コンピューティング機能 2.1 デバイスの倍精度スループットは、1 サイクルあたり 4 回 (DFMA を実行する場合は 8 回) です。これは、ディスパッチされたワープで 32 個のスレッドすべてがアクティブであると想定しています。

4 オペレーション/サイクル/SM * 4 SM * 1800 MHz * 2 オペレーション/DFMA = 56 GFLOPS double

計算は、ワープ内のすべてのスレッドがアクティブであることを前提としています。

質問のコードには、DFMA に融合できる 2 つの依存操作が含まれています。cuobjdump -sass を使用して、アセンブリを調べます。同じ SM で複数のワープを起動すると、テストはレイテンシではなく、依存する命令スループットの測定値になります。

于 2013-02-02T04:46:16.657 に答える
3

カーネルの設計には問題があることに注意する必要があります。つまり、このコードを使用して行う測定は、倍精度の命令スループットとはまったく関係がありません。

すべての倍精度演算を含む計算ループの結果はメモリ書き込みで使用されないため、コンパイラの最適化によって削除されます。CUDA 5コンパイラは、カーネルに対して次のPTXを発行します。

.visible .entry _Z12bench_singlePf(
    .param .u32 _Z12bench_singlePf_param_0
)
{
    .local .align 8 .b8     __local_depot0[8];
    .reg .b32   %SP;
    .reg .b32   %SPL;
    .reg .s32   %r<16>;


    mov.u32     %SPL, __local_depot0;
    cvta.local.u32  %SP, %SPL;
    add.u32     %r3, %SP, 0;
    .loc 2 13 1
    cvta.to.local.u32   %r4, %r3;
    // inline asm
    mov.u32     %r1, %clock;
    // inline asm
    // inline asm
    mov.u32     %r2, %clock;
    // inline asm
    st.local.v2.u32     [%r4], {%r2, %r1};
    cvta.const.u32  %r5, $str;
    // Callseq Start 0
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r5;
    .param .b32 param1;
    st.param.b32    [param1+0], %r3;
    .param .b32 retval0;
    .loc 2 13 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r6, [retval0+0];
    }
    // Callseq End 0
    .loc 2 14 1
    sub.s32     %r7, %r2, %r1;
    cvta.const.u32  %r8, $str1;
    st.local.u32    [%r4], %r7;
    // Callseq Start 1
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r8;
    .param .b32 param1;
    st.param.b32    [param1+0], %r3;
    .param .b32 retval0;
    .loc 2 14 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r9, [retval0+0];
    }
    // Callseq End 1
    .loc 2 15 2
    ret;
}

2つのクロックロード命令は隣接しており、他のコードはへの呼び出しのみprintfです。そのPTXには計算ループはありません。

コンパイラがループ結果が未使用であると推測して最適化できないように、カーネルを再設計する必要があります。1つのアプローチは次のようになります。

__global__ 
void bench_single(float *data, int flag=0)
{
    int i;
    double x = 1.;
    clock_t start, end;
    start = clock();
    for(i=0; i<1000000; i++) {
        x = x * 2.388415813 + 1.253314137;
    }
    end = clock();
    printf("End and start %d - %d\n", end, start);
    printf("Finished in %d cycles\n", end-start);

    if (flag) {
        data[blockIdx.x] = x;
    }
}

カーネルの最後の条件付き書き込みは、コンパイラーがループを最適化するのを妨げるため、コンパイラーは次のPTXを発行します。

.visible .entry _Z12bench_singlePfi(
    .param .u32 _Z12bench_singlePfi_param_0,
    .param .u32 _Z12bench_singlePfi_param_1
)
{
    .local .align 8 .b8     __local_depot0[8];
    .reg .b32   %SP;
    .reg .b32   %SPL;
    .reg .pred  %p<3>;
    .reg .f32   %f<2>;
    .reg .s32   %r<28>;
    .reg .f64   %fd<44>;


    mov.u32     %SPL, __local_depot0;
    cvta.local.u32  %SP, %SPL;
    ld.param.u32    %r6, [_Z12bench_singlePfi_param_0];
    ld.param.u32    %r7, [_Z12bench_singlePfi_param_1];
    add.u32     %r10, %SP, 0;
    .loc 2 13 1
    cvta.to.local.u32   %r1, %r10;
    // inline asm
    mov.u32     %r8, %clock;
    // inline asm
    mov.f64     %fd43, 0d3FF0000000000000;
    mov.u32     %r27, 1000000;

BB0_1:
    .loc 2 10 1
    fma.rn.f64  %fd4, %fd43, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd5, %fd4, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd6, %fd5, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd7, %fd6, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd8, %fd7, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd9, %fd8, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd10, %fd9, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd11, %fd10, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd12, %fd11, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd13, %fd12, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd14, %fd13, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd15, %fd14, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd16, %fd15, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd17, %fd16, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd18, %fd17, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd19, %fd18, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd20, %fd19, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd21, %fd20, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd22, %fd21, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd23, %fd22, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd24, %fd23, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd25, %fd24, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd26, %fd25, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd27, %fd26, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd28, %fd27, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd29, %fd28, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd30, %fd29, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd31, %fd30, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd32, %fd31, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd33, %fd32, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd34, %fd33, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd35, %fd34, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd36, %fd35, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd37, %fd36, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd38, %fd37, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd39, %fd38, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd40, %fd39, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd41, %fd40, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd42, %fd41, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    fma.rn.f64  %fd43, %fd42, 0d40031B79BFF0AC8C, 0d3FF40D931FE078AF;
    .loc 2 9 1
    add.s32     %r27, %r27, -40;
    setp.ne.s32     %p1, %r27, 0;
    @%p1 bra    BB0_1;

    cvta.to.global.u32  %r5, %r6;
    // inline asm
    mov.u32     %r11, %clock;
    // inline asm
    .loc 2 13 1
    st.local.v2.u32     [%r1], {%r11, %r8};
    cvta.const.u32  %r12, $str;
    // Callseq Start 0
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r12;
    .param .b32 param1;
    st.param.b32    [param1+0], %r10;
    .param .b32 retval0;
    .loc 2 13 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r14, [retval0+0];
    }
    // Callseq End 0
    .loc 2 14 1
    sub.s32     %r15, %r11, %r8;
    cvta.const.u32  %r16, $str1;
    st.local.u32    [%r1], %r15;
    // Callseq Start 1
    {
    .reg .b32 temp_param_reg;
    .param .b32 param0;
    st.param.b32    [param0+0], %r16;
    .param .b32 param1;
    st.param.b32    [param1+0], %r10;
    .param .b32 retval0;
    .loc 2 14 1
    call.uni (retval0), 
    vprintf, 
    (
    param0, 
    param1
    );
    ld.param.b32    %r17, [retval0+0];
    }
    // Callseq End 1
    .loc 2 16 1
    setp.eq.s32     %p2, %r7, 0;
    @%p2 bra    BB0_4;

    .loc 2 17 1
    cvt.rn.f32.f64  %f1, %fd43;
    mov.u32     %r18, %ctaid.x;
    shl.b32     %r19, %r18, 2;
    add.s32     %r20, %r5, %r19;
    st.global.f32   [%r20], %f1;

BB0_4:
    .loc 2 19 2
    ret;
}

ここで、コンパイラがループを部分的に展開した場所から、浮動小数点の乗算追加命令の素晴らしいストリームがあることに注意してください。

グレッグ・スミスが指摘したように、命令スケジューリングのレイテンシーを克服するのに十分なワープが特定のSMで実行されるまで、命令スループットの実際の測定値を取得することを期待するべきではありません。これはおそらく、少なくとも1つの大きなブロックを実行する必要があることを意味します。また、printf呼び出しは、スループットに大きな悪影響を与えることに注意してください。ブロックごとにスレッドが1つしかない場合は、結果を書き出すか、（さらに良いことに）グローバルメモリに保存すると、より代表的な番号が得られます。多数のブロックを実行すると、平均化できる多数の測定値が得られます。最後のチェックとして、オブジェクトコードを次のように逆アセンブルする必要があります。cudaobjdumpアセンブラがクロック読み取り命令の位置を移動しないようにするため。そうしないと、依存しているカーネル内のタイミングが壊れます。古いバージョンのアセンブラには、カーネルCコードまたはPTXに挿入された一連のクロック読み取りの機能を壊す可能性のある命令の並べ替えの習慣がありました。

于 2013-02-02T07:52:35.160 に答える