0

この質問には詳細がありません。そこで、この質問を編集する代わりに、別の質問を作成することにしました。新しい質問は次のとおりです。コードを並列化できますか、それとも価値がありませんか?

CUDA で実行しているプログラムがあり、コードの一部がループ内で実行されています (以下に示すようにシリアル化されています)。このコードは、アドレスや NULL ポインターを含む配列内の検索です。すべてのスレッドが以下のコードを実行します。

while (i < n) {
    if (array[i] != NULL) {
        return array[i];
    }
    i++;
}
return NULL;

nのサイズと配列arrayは共有メモリにあります。NULL (最初の一致) とは異なる最初のアドレスにのみ関心があります。

コード全体 (コードの一部だけを投稿しました。コード全体は大きい) は高速で実行されますが、コードの「心臓部」(つまり、より繰り返される部分) は、ご覧のとおりシリアル化されています。この部分（検索）を最適化されたアルゴリズムで並列化できるかどうかを知りたいです。

私が言ったように、プログラムはすでにCUDA（およびデバイスの配列）にあるため、ホストからデバイスへ、またはその逆のメモリ転送はありません。

私の問題は次のとおりnです。大きくありません。8を超えることは難しいでしょう。

並列化を試みましたが、「新しい」コードは上記のコードよりも時間がかかりました。

リダクションと最小演算を勉強していましたが、nが大きい場合に役立つことを確認しました。

それで、何かヒントはありますか？効率的に、つまり低いオーバーヘッドで並列化できますか?

4

2 に答える 2

1

は共有メモリ リソースであると言うのでarray、この検索の結果はブロックの各スレッドで同じです。これは、最初の単純な最適化は、1 つのスレッドのみに検索を実行させることであることを意味します。これにより、ブロックの最初のワープを除くすべての作業が解放されます (結果を待つ必要はありますが、コンピューティング リソースを浪費する必要はありません)。

__shared__ void *result = NULL;
if(tid == 0)
{
    for(unsigned int i=0; i<n; ++i)
    {
        if (array[i] != NULL)
        {
            result = array[i];
            break;
        }
    }
}
__syncthreads();
return result;

次に、従来のブロック内リダクションと同様に、スレッドが並列に検索を実行できるようにします。常に<= 64 であることを保証 できる場合は、これを 1 回のワープで行うことができ、検索中の同期は必要ありません (もちろん、最後の完全な同期は除きます)。n

for(unsigned int i=n/2; i>32; i>>=1)
{
    if(tid < i && !array[tid])
        array[tid] = array[tid+i];
    __syncthreads();
}

if(tid < 32)
{
    if(n > 32 && !array[tid]) array[tid] = array[tid+32];
    if(n > 16 && !array[tid]) array[tid] = array[tid+16];
    if(n > 8 && !array[tid]) array[tid] = array[tid+8];
    if(n > 4 && !array[tid]) array[tid] = array[tid+4];
    if(n > 2 && !array[tid]) array[tid] = array[tid+2];
    if(n > 1 && !array[tid]) array[tid] = array[tid+1];
}

__syncthreads();    
return array[0];

もちろん、この例nでは 2 の累乗を想定しています (それに応じarrayて にNULLs を追加します) が、必要に応じて自由に調整し、これをさらに最適化してください。

于 2013-07-25T12:01:22.467 に答える