0

だから私は非常に基本的な 2D セルオートマトン プログラムを持っています。注意が必要な部分、または少なくとも私が GPU プログラミングを学習するようになった理由は、私の行列が最低でも 1 億要素になる傾向があることです。シリアル プログラムの単純な実装 (第 1 ドラフト) と呼ぶものを CUDA C に作成しました。これはうまく機能しますが、さらに改善する方法があると思います。

テスト目的で、行列の各反復間で生細胞数を追跡する必要があります (私の実装では、生細胞は 1 で表され、死細胞は 0 で表されます)。私の最初のドラフトでは、ホストは、デバイスがマトリックスの更新を終了した後、生細胞をカウントする N^2 のネストされたループを担当します。

より良いアイデアを提供するための疑似コードを次に示します（次のコード自体は、任意の回数の反復後に終了するループ内にあります）。

cuda_kern<<<dimGrid, dimBlock>>>( d_grid );
cudaMemcpy( h_grid, d_grid, matrixSize*sizeof(int), cudaMemcpyDeviceToHost );

*liveCell = 0;
for(a=0; a<gridLength; a++)
    for(b=0; b<gridLength; b++)
        if(h_grid[a][b]==1) (*liveCell)++;

printf("live cell count is %d\n", *liveCell);

ご覧のとおり、かなりコストのかかる操作が 2 つあります。まず、反復ごとにデバイスとホストの間でマトリックス (上記のグリッドと呼ばれる) を転送する必要があります。次に、生細胞数のマトリックスを組み合わせ/カウントする N^2 操作があります。各反復でホストとデバイス間で生細胞数のみを転送することで、これらの両方の手順を排除できると思います。実際、私はこれをコード化することに非常に近づいていますが、質問のタイトルが示すように、シリアル プログラムと最新の CUDA プログラムから確認されたライブ カウントからわずかな「ランダムな」偏差があることを示しています。

これは、私の第 2 ドラフト カーネルの疑似コードです。

__global__ void cuda_kern( liveCell, grid )
{
    i = threadIdx.x + blockIdx.x * blockDim.x
    j = threadIdx.y + blockIdx.y * blockDim.y

    if( i < gridLength AND j < gridLength )
    {
        __shared__ int temp[number of threads in block] 
        x, y, count, changeCounter, sum = 0

        // nested for loops here that update int variable 'count'
        // for those familiar with CA it is a basic neighborhood analysis
        for x to arbitrary neighborhood range
            for y to arbitrary neighborhood range
            {
                //count neighbors, that is update 'count' variable
            }

        // __syncthreads()

        if( grid[i*gridLength+j] == 0 AND count == 5 ) 
        { // NOTE: 'count' above could be any arbitrary integer 1 - 8  
            grid[i*gridLength+j] = 1
            changeCounter += 1
        }
        else if( grid[i*gridLength+j] ==  1 AND count >= 5 )
        {
            grid[i*gridLength+j] = 0
        }
        else 
        {
            grid[i*gridLength+j] = 1
            changeCounter += 1
        }

        //__syncthreads()

        temp[threadIdx.x] = changeCounter 

        //__syncthreads()

        if (threadIdx.x == 0)
            for( i = 0; i < N; i++ ) 
            { 
                sum += temp[i]
                //__syncthreads()
            }

        atomicAdd( liveCell, sum )

    } // end of if ( i < gridLength and j < gridLength )   
    return;
} // end of kernel

上記の疑似コードの説明: セル オートマトン関数の背後にあるロジックの大部分など、本質的でないものを除外しようとしました。私のプログラムでは変数はすべて整数で、行列自体は整数型です。スレッドを同期しようとした場所をコメントに残しました（役に立たなかった）。

私のロジックの説明: 最初のドラフトの実装について説明したときに前述したように、デバイスとホストの間でマトリックスを転送するコストのかかる操作と、反復ごとにホスト上の生細胞をカウントする N^2 コストを回避したいと考えています。これを行う鍵は、変数「changeCounter」で確認できます。基本的に、「changeCounter」は、セルが「1」または生きていると指定されるたびに 1 ずつ増加する必要があります。「changeCounter」の値を格納するために、現在のスレッドによってインデックス付けされている共有変数「temp」を使用しようとしています。ブロック内のすべてのスレッドが終了したら、'temp' 配列を 1 つの変数 'sum' に圧縮しようとします。これをアトミック操作 'atomicAdd' を介して liveCell に追加します。

私が抱えている問題は、この方法での結果が（より高速ですが）元のドラフトの結果と一致しないことです。また、同じマトリックスでプログラムを実行するたびに、ライブセルの反復ごとにわずかな偏差がありますカウント。

UPDATE 1k x 1k マトリックスの出力例:

第 1 ドラフトの GPU 実装では、毎回次の生細胞数が生成されます。これらは、2 番目のドラフトで作成したい結果です。

initial live cell count: 393592 
itr. 0  live cell count: 364118
itr. 1  live cell count: 315417
itr. 2  live cell count: 300413 
itr. 3  live cell count: 284503

2nd Draft GPU 実装のライブ セル カウントは毎回わずかに異なりますが、上記に近い値です。ここでは、アイデアを提供するために 3 つの個別の実行を示します。

実行 A:

initial live cell count: 393592
itr. 0  live cell count: 372402
itr. 1  live cell count: 324114
itr. 2  live cell count: 309580
itr. 3  live cell count: 291393

実行 B:

initial live cell count: 393592
itr. 0  live cell count: 374139
itr. 1  live cell count: 323948
itr. 2  live cell count: 307214
itr. 3  live cell count: 292582

実行 C:

initial live cell count: 393592
itr. 0  live cell count: 372391
itr. 1  live cell count: 323105
itr. 2  live cell count: 308295
itr. 3  live cell count: 292512

最初の生細胞数は毎回ホストで計算されます。これが、すべての例で一貫している理由を説明しています。

更新の終了

お時間をいただきありがとうございます。さらに情報が必要な場合は、お知らせください。

4

1 に答える 1

1

これを行うべきではありません：

    if (threadIdx.x == 0) {
        for( i = 0; i < N; i++ ) 
        { 
            sum += temp[i]
            //__syncthreads()
        }
        atomicAdd( liveCell, sum )// could probably just update the memory without atomic
    }

ブロック内のすべてのスレッドが合計を計算すると言ったように、スレッド 0 ですべてのスレッドの合計を合計し、それを指定されたポインタに追加する必要があります。すべてのスレッドがatomicAddを実行する場合、それは異なります。おそらくスレッド0でのみ実行する必要があり、おそらくアトミック操作を削除できます。

また、そのループで削減を行います。

for (i=blockDim.x/2; i>0; i>>=1) {
     if(threadIdx.x < i) {
         temp[threadIdx.x] += temp[threadIdx.x + i];
     }
    __syncthreads();
}
if(threadIdx.x == 0) {
    sum = temp[0];
    liveCell += sum;
}
于 2012-12-27T22:54:37.207 に答える