1

2 つの配列の合計二乗誤差を計算する次の CUDA カーネル コードがあります。

__global__ void kSquaredError(double* data, double* recon, double* error, 
                               unsigned int num_elements)
{
    const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    for (unsigned int i = idx; i < num_elements; i += blockDim.x * gridDim.x) {
        *error += pow(data[i] - recon[i], 2);
    }
}

単一のスカラー出力が必要です (エラー)。この場合、すべてのスレッドが同時にエラーに書き込んでいるように見えます。同期する必要がある方法はありますか？

現在、私は悪い結果を得ているので、何か問題があると推測しています。

4

1 に答える 1

1

現在行っている実装は、すべてのスレッドが同じグローバル メモリ アドレスを同時に更新しようとするため、競合状態の影響を受けます。atomicAdd代わりに関数を簡単に配置できますが、*error += pow...更新ごとにシリアル化されるため、パフォーマンスの問題が発生します。

代わりに、次のように、共有メモリを使用して削減を試みてください。

_global__ void kSquaredError(double* data, double* recon, double* error, unsigned int num_elements) {
    const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int tid = threadIdx.x;
    __shared__ double serror[blockDim.x];//temporary storage of each threads error

    for (unsigned int i = idx; i < num_elements; i += blockDim.x * gridDim.x) {
        serror[tid] += pow(data[i] - recon[i], 2);//put each threads value in shared memory
    }

    __syncthreads();

    int i = blockDim.x >> 1; //halve the threads
    for(;i>0;i>>=1) {//reduction in shared memory
            if(tid<i) {
                serror[tid] += serror[tid+i];
                __syncthreads();//make shure all threads are at the same state and have written to shared memory
            }
    }

    if(tid == 0) {//thread 0 updates the value in global memory
        atomicAdd(error,serror[tid]);// same as *error += serror[tid]; but atomic
    }
}

次の原則に従って動作します。各スレッドには独自の一時変数があり、すべての入力のエラーの合計を計算し、すべてのスレッドが__syncthreads命令に収束してすべてのデータが完全であることを確認します。

ブロック内のすべてのスレッドの半分は、対応する残りの半分から1つの値を取得し、それを独自の値に追加し、半分のスレッドを再度実行して、合計が1つのスレッド(スレッド0)になるまで繰り返します。 .

スレッド 0 は、atomicAdd 関数を使用してグローバル メモリを更新し、他のブロックがある場合に競合状態を回避します。

最初の例を使用して、すべての割り当てでatomicAddを使用するとします。シリアル化されるアトミック関数がありましたgridDim.x*blockDim.x*num_elementsが、今ではgridDim.xアトミック関数しかなく、はるかに少ないです。

cuda を使用したリダクションの仕組みについて詳しくは、CUDA での並列リダクションの最適化を参照してください。

編集

リダクション for ループに if ステートメントを追加しましたが、それを忘れていました。

于 2013-01-21T20:56:58.933 に答える