1

最近、CUDA を使用してリダクションのアルゴリズムをテストしました (たとえば、http: //www.cuvilib.com/Reduction.pdf の 16 ページにあります)。しかし、最終的には原子性を使用しないという問題に遭遇しました。したがって、基本的には各ブロックの合計を計算して共有配列に格納します。次に、グローバル配列 x に戻します (tdx は threadIndex.x で、i はグローバル インデックスです)。

if(i==0){
        *sum = 0.; // Initialize to 0
    }
__syncthreads();
if (tdx == 0){       
    x[blockIdx.x] = s_x[tdx]; //get the shared sums in global memory
}
__syncthreads();

次に、最初の x 要素 (ブロックの数だけ) を合計します。アトミック性を使用すると正常に動作しますが (CPU と同じ結果)、以下のコメント行を使用すると動作せず、多くの場合「nan」が生成されます。

if(i == 0){    
    for(int k = 0; k < gridDim.x; k++){
        atomicAdd(sum, x[k]); //Works good
       //sum[0] += x[k]; //or *sum += x[k]; //Does not work, often results in nan
    }
}

実際、atomicadd を直接使用して共有合計を合計していますが、これが機能しない理由を理解したいと思います。操作を単一のスレッドに制限する場合、アトミックな追加はまったくナンセンスです。そして、単純な合計はうまくいくはずです!

4

2 に答える 2

4

__syncthreads()異なるブロック間ではなく、同じブロック内のスレッドのみを同期し、CUDA にはブロック間の安全な同期メカニズムがありません。

誤った結果は、同期の問題が原因です。オペランドx[k]は、さまざまなブロックからの計算の結果です。x[0]は blockの結果0、x[1]はblock の結果1などです。一部のブロックが実際に計算を終了する前に、スレッド0がそれらの加算を開始する可能性があります。

2 番目のコード スニペットを別のカーネルに配置する必要があります。これにより、同期が強制され、ラインsum[0] += x[k];が機能するようになります。

于 2013-07-24T10:30:57.237 に答える
0

指摘されているように、問題はブロック間で同期できないため、最初のパスの後に同期が欠落していることが原因です。ツールキットで提供されるサンプル コードには、リダクションに関する適切なウォークスルーがあります。

そうは言っても、リダクション カーネル (またはスキャンなどの他のプリミティブ) がライブラリ コードに存在する場合は、このようなプリミティブを作成しないことを強くお勧めします。他の場所に労力を注ぎ、利用可能な既存の最適化されたコードを再利用することをお勧めします。もちろん、これを学ぶためにこれを行っている場合、これは当てはまりません!

ThrustとCUBをご覧になることをお勧めします。

于 2013-07-24T13:07:57.500 に答える