0

CUDA Fermi GPU で非常に単純なカーネル <<<1,512>>> を起動します。

__global__ void kernel(){
int x1,x2;

x1=5;
x2=1;

for (int k=0;k<=1000000;k++)
  {
   x1+=x2;

  }
}

カーネルは非常に単純で、10^6 回の加算を行い、グローバル メモリには何も転送しません。結果は正しいです。つまり、ループ x1 の後 (512 個のスレッド インスタンスすべてで) には 10^6 + 5 が含まれます。

カーネルの実行時間を測定しようとしています。Visual Studio のパラレル nsight と nvvp の両方を使用します。Nsight は 2.5 マイクロ秒を測定し、nvvp は 4 マイクロ秒を測定します。

問題は次のとおりです。たとえば、ループのサイズを 10^8 に大幅に増やしても、時間は一定のままです。ループサイズを大幅に減らしても同じです。なぜこれが起こるのですか？

ループ内で共有メモリまたはグローバル メモリを使用する場合、測定値は実行中の作業を反映することに注意してください (つまり、比例関係があります)。

4

2 に答える 2

4

前述のように、CUDA コンパイラの最適化はデッド コードの削除に非常に積極的です。x2メモリに書き込まれる値には関与しないため、それとループを取り除くことができます。コンパイラは、コンパイル時に推定できる結果も事前に計算するため、ループ内のすべての定数がコンパイラに認識されている場合、コンパイラは最終結果を計算して定数に置き換えることができます。

これらの問題の両方を回避するには、コードを次のように書き直します。

__global__ 
void kernel(int *out, int x0, bool flag)
{
    int x1 = x0, x2 = 1;

    for (int k=0; k<=1000000; k++) {
       x1+=x2;
    }

    if (flag) out[threadIdx.x + blockIdx.x*blockDim.x] = x1;
}

そして、次のように実行します。

kernel<<<1,512>>>((int *)0, 5, false);

の初期値をx1引数としてカーネルに渡すことで、ループの結果がコンパイラで使用できないようにします。フラグはメモリ ストアを条件付きにし、メモリ ストアは計算全体を安全に削除できないようにします。実行時にフラグが false に設定されている限り、ストアは実行されないため、ループのタイミングには影響しません。

于 2013-08-30T05:02:12.167 に答える