0

CUDA を使用しない場合、私のコードは、システム内のすべての座標ペア間の距離を計算し、それらの距離をビンに並べ替える 2 つの for ループだけです。

私の CUDA バージョンの問題は、明らかに複数のスレッドが同時に同じグローバル メモリ ロケーションに書き込めないことです (競合状態?)。スレッドの 1 つだけが各ビンに書き込みを終了したため、ビンごとに得られる値は正しくありません。

__global__ void computePcf(
        double const * const atoms,
        double * bins,
        int numParticles,
        double dr) {

    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numParticles - 1) {
        for (int j = i + 1; j < numParticles; j++) {
            double r = distance(&atoms[3*i + 0], &atoms[3*j + 0]);

            int binNumber = floor(r/dr);

            // Problem line right here.
            // This memory address is modified by multiple threads
            bins[binNumber] += 2.0;
        }
    }
}

だから…どうすればいいのかわからない。私はグーグルで共有メモリについて読んでいますが、問題は、距離の計算を行うまで、どのメモリ領域にアクセスするのかわからないことです!

VMDと呼ばれるプログラムがGPUを使用してこの計算を高速化するため、これが可能であることはわかっています。任意のヘルプ (またはアイデア) をいただければ幸いです。これを最適化する必要はありません。機能的なだけです。

4

1 に答える 1

1

bins[]いくつありますか？bins[]タイプにする必要がある理由はありdoubleますか？あなたのコードからは明らかではありません。あなたが持っているのは本質的にヒストグラム操作であり、高速な並列ヒストグラム手法を見たいと思うかもしれません。 スラストは興味深いかもしれません。

コードで考慮すべきいくつかの手段があります。

  1. スレッドの特定のグループ (またはビンの計算) が互いに踏み込まないように計算を調整するようにアルゴリズムを再構築する方法があるかどうかを確認してください。これは、おそらくソート距離に基づいて達成される可能性があります。

  2. アトミックを使用する これで問題は解決するはずですが、実行時間の点でコストがかかる可能性があります (ただし、非常に単純なので、試してみることをお勧めします)。

    bins[binNumber] += 2.0;

    このようなもの：

    int * bins,
...
atomicAdd(bins+binNumber, 2);

    binsタイプが の場合でもこれを行うことができますがdouble、もう少し複雑です。で行う方法の例については、ドキュメントを参照してatomicAddくださいdouble。

  3. 数binsが少ない場合 (おそらく数千以下)、複数のスレッドブロックによって更新されるビンのセットをいくつか作成し、リダクション操作 (ビンのセットを要素ごとに一緒に追加する) を使用できます。処理シーケンスの終了。この場合、カーネル コードに追加のループを配置して、それぞれが複数の要素を処理する少数のスレッドまたはスレッド ブロックを使用することを検討してください。これにより、各パーティクルの処理が完了すると、ループは次の処理にジャンプします。粒子を変数に追加gridDim.x*blockDim.xしi、プロセスを繰り返します。各スレッドまたはスレッドブロックには独自のビンのローカル コピーがあるため、他のスレッド アクセスを踏まずにこれを行うことができます。

    たとえば、int 型のビンが 1000 個しか必要ないとします。1000 セットのビンを作成できましたが、それは約 4 メガバイトしか占めません。次に、1000 のスレッドのそれぞれに独自のビン セットを与えることができます。その後、1000 のスレッドのそれぞれに、更新する独自のビン セットがあり、アトミックは必要ありません。これは、他のスレッドに干渉できないためです。各スレッドが複数のパーティクルをループするようにすることで、マシンを効果的にビジー状態に保つことができます。すべての粒子ビニングが完了したら、おそらく別のカーネル呼び出しを使用して、1000 個のビンセットを一緒に追加する必要があります。

于 2013-07-10T01:25:39.620 に答える