1

GPUで静的ステートメントを実行できるかどうか疑問に思いました

これは私のコードの一部です

__global__ void run_state(float *oldv, float* newv, float* w, int t)
{
  int i = threadIdx.x;

  nextState(i, oldv, newv, w, t);


  newv[0]   = newv[1];
  newv[N+1] = newv[N];

}

ループして一度だけ実行した後newv[0] = newv[1];にnewv[N+1] = newv[N];実行できるかどうか疑問に思って いましたか？nextState

4

2 に答える 2

2

あなたが何を達成したいのか正確にはわかりません（何ごとに1回実行しますか？）ので、いくつかの代替案を示します。

ブロックごとに 1回ステートメントを実行する場合は、スレッド インデックスをテストするだけで実行できます。

__syncthreads();
if ((threadIdx.x | threadIdx.y | threadIdx.z) == 0) {
    // statements that are only executed once per block
}
__syncthreads();

カーネル呼び出しごとにステートメントを 1 回実行する場合は、カーネルのブロックが特定の順序で実行されないため、ステートメントをいつ実行するかをもう少し具体的にする必要があります。

上記の例を単純に拡張すると、ステートメントがカーネル呼び出しごとに 1 回実行されるバージョンが得られますが、時間は指定されていません。

if ((threadIdx.x | threadIdx.y | threadIdx.z
     | blockIdx.x | blockIdx.y | blockIdx.z) == 0) {
    // statements that are executed once per kernel invocation,
    // at an unspecified time
}

追加のグローバル メモリ帯域幅の使用を犠牲にして、グローバル メモリでアトミック操作を使用してステートメントを実行するタイミングを指定することができます。パフォーマンスへの影響を制限するには、通常、これらのアトミック操作をブロックごとに 1 回だけ実行することをお勧めします。

最初のブロックがそれに到達するとすぐにステートメントを実行したい場合は、グローバル フラグの簡単なテストで十分です。

__global__ volatile unsigned int statementHasExecuted;

...

__syncthreads();
if ((threadIdx.x | threadIdx.y | threadIdx.z) == 0) {
    unsigned int state = atomicMin((unsigned int*)&statementHasExecuted, 1);
    if (state == 0) {
        // statements that are executed once per kernel invocation,
        // as soon as the first block reaches the statement

        // now make the results visible to the other blocks:
        __threadfence();
        // and signal that we are done:
        atomicMin((unsigned int*)&statementHasExecuted, 2);
    } else {
        while (state < 2) {
            // busy wait until execution of statement in other block has finished:
            state = statementHasExecuted;
        }
    }
}
__syncthreads();

statementHasExecutedカーネルを起動するたびにゼロにリセットする必要があります。

あなたが求めているのは、他のすべてのブロックが終了した後にのみ実行されるコードの一部である、逆のケースかもしれません。プログラミング ガイドの付録 B.5「メモリ フェンス関数」には、このためのサンプル コードが示されています。

(すべてのコードはブラウザーで作成され、テストされていないことに注意してください。エラーが含まれている可能性があります)。

于 2012-12-06T15:43:57.960 に答える
0

最終的な作業を行うスレッドを 1 つ指定します。例えば：

__global__ void run_state(float *oldv, float* newv, float* w, int t)
{
  int i = threadIdx.x;

  nextState(i, oldv, newv, w, t);

  if (i == 0) {
    newv[0]   = newv[1];
    newv[N+1] = newv[N];
  }
}
于 2012-12-06T15:16:40.193 に答える