4

GTX680の実際のパフォーマンスに疑問があるため、GTX680の計算パフォーマンスをテストしようとしています。同じ結果が得られたら、誰かが彼のGTX 680でテストできるのではないか、あるいはカードからより多くのパフォーマンスを引き出すために何ができるのか教えてくれるのではないかと思っていました。

私はこの小さなプログラムを書きました

#include <stdlib.h>
#include <stdio.h>
#include <float.h>
#include "cuComplex.h"
#include "time.h"
#include "cuda_runtime.h"
#include <iostream>
using namespace std;
__global__ void test(int loop, int *out)
{
    register int a=0;
    for (int x=0;x<loop;x++)
    {
        a+=x*loop;
    }


    if (out!=NULL) *out=a;


}
int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaDeviceReset();
    cudaDeviceReset();
    cudaSetDevice(0);
    cudaThreadSetCacheConfig(cudaFuncCachePreferShared);        
    // Allocate and generate buffers
    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    int b=1000; 
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    test<<<blocks,threadsPerBlock,0>>>(300,
            NULL
            );

    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f", timestamp);
}

nvccでコンパイルするとこのPTXが得られます

//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sat Sep 22 02:35:14 2012 (1348274114)
// Cuda compilation tools, release 5.0, V0.2.1221
//

.version 3.1
.target sm_30
.address_size 64

    .file   1 "/tmp/tmpxft_00000e7b_00000000-9_perf.cpp3.i"
    .file   2 "/opt/home/daniel/a/perf.cu"

 .visible .entry _Z4testiPi(
    .param .u32 _Z4testiPi_param_0,
    .param .u64 _Z4testiPi_param_1
 )
 {
    .reg .pred      %p<4>;
    .reg .s32       %r<15>;
    .reg .s64       %rd<3>;


    ld.param.u32    %r6, [_Z4testiPi_param_0];
    ld.param.u64    %rd2, [_Z4testiPi_param_1];
    cvta.to.global.u64      %rd1, %rd2;
    mov.u32         %r13, 0;
    .loc 2 12 1
    setp.lt.s32     %p1, %r6, 1;
    mov.u32         %r14, %r13;
    mov.u32         %r11, %r13;
    @%p1 bra        BB0_2;

 BB0_1:
    .loc 2 14 1
    mad.lo.s32      %r14, %r11, %r6, %r14;
    .loc 2 12 20
    add.s32         %r11, %r11, 1;
    .loc 2 12 1
    setp.lt.s32     %p2, %r11, %r6;
    mov.u32         %r13, %r14;
    @%p2 bra        BB0_1;

 BB0_2:
    .loc 2 18 1
    setp.eq.s64     %p3, %rd2, 0;
    @%p3 bra        BB0_4;

    .loc 2 18 1
    st.global.u32   [%rd1], %r13;

 BB0_4:
    .loc 2 21 2
    ret; 
 }

カーネルは1.936msで実行されます

私の計算によると、GFLOPSのパフォーマンスは3TFLOPSの理論値のわずか3分の1の1.1 TFLOPでした（参照： http： //www.geforce.com/hardware/desktop-gpus/geforce-gtx-680）。なぜこれほど遅いのですか。 ？

私の計算の詳細は次のとおりです

mad.lo.s32      %r14, %r11, %r6, %r14;  //2 FLOPS
.loc 2 12 20
 add.s32         %r11, %r11, 1;     //1 FLOP
.loc 2 12 1
 setp.lt.s32     %p2, %r11, %r6;    //1 FLOP
 mov.u32         %r13, %r14;        // 1 FLOP
 @%p2 bra        BB0_1;             //1 FLOP

 + 1 FLOP (just as a buffer as I don't know branching how much it takes)

ループ内の1回の反復の合計FLOPSは7FLOPSです

反復のみを考慮する

スレッドごとに300回の反復があります1024*1000ブロックがあります

総反復回数FLOPS=300 * 1024 * 1000 * 7 = 2.15 GFLOPS

総カーネル時間は1.936msです

したがって、スループット= 1.11 TFLOPS

助けてくれてありがとう

ダニエル

4

4 に答える 4

3

このサンプルプログラムは、@RobertCrovellaの回答に基づいています。Robertのカーネルは、データの依存関係によって制限されています。FMA命令間のデータ依存関係を減らすことにより、このカーネルはGTX680で2.4〜2.5TFLOPSを達成する必要があります。

現在の実装では、命令フェッチとデータ依存関係が制限されています。カーネルを微調整して、達成されたFLOPSをさらに10％改善できる必要があります。

Nsight Visual Studio Edition2.xおよび新しい3.0RC候補は、このカーネルを分析するために必要なメトリックを提供します。

2.xおよび3.0では、次の実験を使用してカーネルを分析する必要があります。

  1. 命令統計-SMアクティビティ-すべてのSMが100％に近いことを確認します
  2. 発行効率-適格ワープ-Keplerの場合各ワープスケジューラがサイクルごとに命令を発行するには、アクティブサイクルあたりの適格ワープが4より大きくなければなりません。
  3. 発行効率-発行ストール-ワープ発行効率は、適格なワープの数が不十分なために各ワープスケジューラーが発行できなかった頻度を指定します。これが高い場合、IssueStallReasonsはリミッターを特定するのに役立ちます。
  4. 達成されたFLOP-これは、カーネルによって実行される単精度および倍精度浮動小数点演算のタイプとレートの両方の内訳を表示します。

Robertのカーネルの場合、各命令には書き込み後の読み取り依存関係があるため、実行依存関係は非常に高くなりました。命令レベルの並列性を高めることで、パフォーマンスが3倍になりました。現在、カーネルは主に命令フェッチに制限されています。

新しいNsightVSE3.0 RC（現在入手可能）には、実行された命令の数や命令ごとのアクティブなスレッドの数など、命令ごとの統計で注釈が付けられたアセンブリまたはソースコードも表示されます。この例では、ツールを使用してデータの依存関係を特定し、コンパイラが理論的に達成されたFLOPSの50％を超える値に到達するために必要なFMA命令を生成していることを確認できます。

__global__ void test(float loop, float *out)
{
    register float a=1.0f;
    register float b=1.0f;
    register float c=1.0f;
    register float d=1.0f;
    register float e=1.0f;
    register float f=1.0f;
    register float g=1.0f;
    register float h=1.0f;

    for (float x=0;x<loop;x++)
    {
        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;

        a+=x*loop;
        b+=x*loop;
        c+=x*loop;
        d+=x*loop;
        e+=x*loop;
        f+=x*loop;
        g+=x*loop;
        h+=x*loop;
    }
    if (out!=NULL) *out=a+b+c+d+e+f+g+h;
}

int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaDeviceReset();
    cudaSetDevice(0);
    cudaThreadSetCacheConfig(cudaFuncCachePreferShared);
    // Allocate and generate buffers
    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    test<<<blocks,threadsPerBlock,0>>>(30,NULL);

    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f\n", timestamp);
}
于 2012-10-30T03:10:14.623 に答える
1

問題は、整数の乗算を使用していることだと思います。コンピューティング機能3.0アーキテクチャでの32ビット整数乗算は、32ビット浮動小数点スループットのわずか1/6です（CUDA Cプログラミングガイドバージョン5.5から抜粋した以下の表を参照してください）。3.0アーキテクチャの32ビット整数乗算パフォーマンスと32ビット浮動小数点パフォーマンスを比較します。

コンピューティングアプリケーションで主に使用される他のいくつかの整数演算と型変換では、3.0で同様にパフォーマンスが低下します。

ここに画像の説明を入力してください

于 2013-11-13T06:51:06.357 に答える
0

このコードでより良い結果が得られるかどうかを確認してください。これは単なる例であり、コードとまったく同じことを行うわけではなく、フロップを再カウントする必要があると思います。

#include <stdio.h>
using namespace std;
__global__ void test(float loop, float *out)
{
    register float a=1.0f;
    for (float x=0;x<loop;x++)
    {
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
        a+=x*loop;
    }


    if (out!=NULL) *out=a;


}
int main(int argc, char *argv[])
{
    float timestamp;
    cudaEvent_t event_start,event_stop;
    // Initialise
    cudaDeviceReset();
    cudaSetDevice(0);
    cudaThreadSetCacheConfig(cudaFuncCachePreferShared);
    // Allocate and generate buffers
    cudaEventCreate(&event_start);
    cudaEventCreate(&event_stop);
    cudaEventRecord(event_start, 0);
    dim3 threadsPerBlock;
    dim3 blocks;
    threadsPerBlock.x=32;
    threadsPerBlock.y=32;
    threadsPerBlock.z=1;
    blocks.x=1;
    blocks.y=1000;
    blocks.z=1;

    test<<<blocks,threadsPerBlock,0>>>(30,
            NULL
            );

    cudaEventRecord(event_stop, 0);
    cudaEventSynchronize(event_stop);
    cudaEventElapsedTime(&timestamp, event_start, event_stop);
    printf("Calculated in %f\n", timestamp);
}

これをarch=sm_20またはsm_30でコンパイルすると、カーネルループ内で10個のfma命令が連続して取得され、コードは介在しません。コードよりも高速で、理論上のフロップのピークに近いと思います。はい、整数OP/秒と浮動小数点OP/秒には違いがあります。このコードを実行する場合は、コメントして、計算されたパフォーマンスを教えてください。

于 2012-10-29T16:11:50.513 に答える
0

テストカーネルは、浮動小数点演算ではなく整数演算を実行しています。したがって、FLOPSは、そのカーネルにとってすべて一緒に間違ったメトリックです。

FLOPS = FLoating point Operations Per Second

ただし、元の質問に戻ると、GPUは整数計算ではなく浮動小数点計算用に最適化されているため、カーネルは低速です。

適切なテストを行うには、整数ではなく浮動小数点数を使用するようにテストカーネルを変換してみてください。

さらに、FLOPSをステップに注釈を付けるループでは、FLOPSは1秒あたりのメジャーであり、整数演算であるため、やはり意味がありません。変換したら、1秒あたりの浮動小数点演算ではなく、個々の浮動小数点演算としてカウントします。

于 2012-10-29T16:23:44.747 に答える