0

間にカウンターがあるネストされたループがあります。外側のループに CUDA インデックスを使用することはできましたが、この種のループでさらに並列処理を利用する方法が思い浮かびません。それに似たものを扱った経験はありますか？

int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < Nx) {
    counter = 0;
    for (k = 0; k < Ny; k++) {

        d_V[i*Ny + k] = 0;

        if ( d_X[i*Ny + k] >= 2e2 ) {

             /* do stuff with i and k and counter i.e.*/
                d_example[i*length + counter] = k;
                    ...
             /* increment counter */
             counter++;
        }
    }
}

私が見ている問題は、カウンターを処理する方法ですk.CUDAでインデックスを作成することもできますthreadIdx.y + blockIdx.y * blockDim.y

4

3 に答える 3

1

可能であれば、cudppまたはthrust（remove_ifやcompactなどの並列関数を実装するライブラリ-何か、例にあるもの）を使用できます。

カドップ

スラスト

これらのページでは、簡単な例とその使用方法を見つけることができます。IMHOは推力よりも速いので、私はcudppを好みます。

于 2012-10-01T13:48:13.427 に答える
1

ループ反復間で使用されるカウンター/ループ変数を持つことは、並列化に対する自然なアンチテーゼです。理想的な並列ループには、相互の知識がなくても、任意の順序で実行できる反復があります。残念ながら、共通の変数は、順序に依存し、相互に認識します。

カウンターを使用して、d_example配列を隙間なくパックしているようです。この種のことは、いくらかのメモリを浪費することにより、計算時間をより効率的にすることができます。設定されない d_example の要素をゼロのままにしておくと、 を非効率的にパッキングd_exampleして、高価な計算ステップの後で d_example でフィルターを実行できます。

実際、配列が読み取られるときにフィルター処理を変更されたイテレーターに任せることもできます。これはゼロ値をスキップするだけです。ゼロが配列内の有効な値である場合は、特定の NaN 値または別のマスク配列を使用してください。

int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < Nx) {
    for (k = 0; k < Ny; k++) {

        d_V[i*Ny + k] = 0;

        if ( d_X[i*Ny + k] >= 2e2 ) {

             /* do stuff with i and k and counter i.e.*/
                d_example[i*length + i*k] = k;
                d_examask[i*length + i*k] = 1;
                    ...
             /* increment counter */
        } else {
             d_examask[i*length+i*k] = 0;
        }
    }
}
于 2012-10-01T11:03:39.110 に答える
1

配列の 2 番目のインデックスとして threadIDx.y を使用できることに注意してください。詳細については、こちらを参照してください: http://www.cs.sunysb.edu/~mueller/teaching/cse591_GPU/threads.pdf

たとえば、2 次元のブロックがある場合、threadix.x と threadix.y をインデックスとして使用し、ワークグループのオフセット (blockidx.x * blockDim.x) をオフセットとして追加できます。

分岐は GPU では非常にコストがかかり、特定のワークグループ内のすべてのスレッドは、グループ内のすべてのタスクが続行されるのを常に待機するため、単純にすべての要素を計算し、必要のない要素を破棄することをお勧めします。カウンターの使用を完全に回避できる可能性があります。そうでない場合、最善の解決策は、彼のコメントで phoad によって指定されているように、グローバル カウンターで CUDA API のアトミック インクリメント機能を使用することです。

于 2012-10-01T11:18:34.340 に答える