4

私はいくつかの比較を行い、2つのオブジェクトが衝突するかどうかを決定するカーネルを持っています。衝突するオブジェクトのIDを出力バッファに保存したい。出力バッファにギャップを持たせたくありません。各衝突を出力バッファの一意のインデックスに記録したいと思います。

そこで、共有メモリ（ローカル合計）とグローバルメモリ（グローバル合計）にアトミック変数を作成しました。以下のコードは、衝突が見つかったときの共有変数の増分を示しています。今のところ、グローバルメモリでアトミック変数をインクリメントすることに問題はありません。

__global__ void mykernel(..., unsigned int *gColCnt) {
    ...

    __shared__ unsigned int sColCnt;
    __shared__ unsigned int sIndex;

    if (threadIdx.x == 0) {
        sColCnt = 0;
    }

    __syncthreads();

    unsigned int index = 0;
    if (colliding)
        index = atomicAdd(&sColCnt, 1); //!!Time Consuming!!

    __syncthreads();

    if (threadIdx.x == 0)
        sIndex = atomicAdd(gColCnt, sColCnt);

    __syncthreads();

    if (sColCnt + sIndex > outputSize) { //output buffer is not enough
        //printf("Exceeds outputsize: %d + %d > %d\n", sColCnt, sIndex, outputSize);
        return;
    }

    if (colliding) {
        output[sIndex + index] = make_uint2(startId, toId);
    }
}

私の問題は、多くのスレッドがアトミック変数をインクリメントしようとすると、シリアル化されることです。prefix-sumのようなものを書く前に、これを効率的に行う方法があるかどうかを尋ねたかったのです。

この1行があるため、カーネルの経過時間は13ミリ秒から44ミリ秒に増加します。

プレフィックスサムのサンプルコードを見つけましたが、NVIDIAのディスカッションボードがダウンしているため、参照されているリンクが失敗します。 https://stackoverflow.com/a/3836944/596547

編集：上記にコードの最後も追加しました。実際、私には階層があります。すべてのコード行の影響を確認するために、すべてのオブジェクトが互いに衝突するシーン、極端な場合、およびオブジェクトがほとんど衝突しない別の極端な場合を設定します。

最後に、共有アトミック変数をグローバル変数（gColCnt）に追加して、衝突の数を外部に通知し、正しいインデックス値を見つけます。ここでは何らかの方法でatomicAddを使用する必要があると思います。

4

2 に答える 2

1

たとえば、並列ストリーム圧縮アルゴリズムの使用を検討してくださいthrust::copy_if。

于 2012-07-22T14:24:24.670 に答える
1

関連するnvidiaブログ記事：http：//devblogs.nvidia.com/parallelforall/gpu-pro-tip-fast-histograms-using-shared-atomics-maxwell/

于 2015-03-21T07:11:18.583 に答える