CUDAのコンウェイのゲームを最適化するには？

Question

私はライフゲームのために、このCUDAカーネルを書いた：私は、エラー/最適化のために

__global__ void gameOfLife(float* returnBuffer, int width, int height) { 
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; 
    float p = tex2D(inputTex, x, y); 
    float neighbors = 0; 
    neighbors += tex2D(inputTex, x+1, y); 
    neighbors += tex2D(inputTex, x-1, y); 
    neighbors += tex2D(inputTex, x, y+1); 
    neighbors += tex2D(inputTex, x, y-1); 
    neighbors += tex2D(inputTex, x+1, y+1); 
    neighbors += tex2D(inputTex, x-1, y-1); 
    neighbors += tex2D(inputTex, x-1, y+1); 
    neighbors += tex2D(inputTex, x+1, y-1); 
    __syncthreads(); 
    float final = 0; 
    if(neighbors < 2) final = 0; 
    else if(neighbors > 3) final = 0; 
    else if(p != 0) final = 1; 
    else if(neighbors == 3) final = 1; 
    __syncthreads(); 
    returnBuffer[x + y*width] = final; 
} 

探しています。 パラレルプログラミングは私には全く新しいものです。正しい方法を知りたいのかどうかはわかりません。

残りは、入力配列からCUDA配列にバインドされた2DテクスチャinputTexまでのmemcpyです。出力はグローバルメモリからホストへとmemcpyされてから処理されます。

ご覧のとおり、スレッドは1つのピクセルを扱います。いくつかのソースがスレッドごとに行を行うことを提案している最速の方法であるかどうかは分かりません。私が正しくNVidiaを理解している場合は、スレッドが多いほど良いと言います。私は実践的な経験を持つ人からこれに関する助言を愛するだろう。

Answer 1

私の2セントです。

多分、マルチプロセッサとGPUメモリ間の通信の待ち時間に縛られるようです。必要なデータがキャッシュにない場合は、それぞれ30〜50クロックティックを実行するコードがあり、少なくとも3回のメモリアクセスを生成し、それぞれ200クロック以上のクロックが必要です。

テクスチャメモリを使用すると、それに対処する良い方法ですが、必ずしも最適な方法ではありません。

少なくともスレッドごとに（水平方向に）4ピクセルずつ行うようにしてください。グローバルメモリは一度に128バイトにアクセスすることができます（128バイト間隔で任意のバイトにアクセスしようとするワープがあれば、追加費用なしでキャッシュライン全体を引き込むことができます）。ワープは32スレッドであるため、各スレッドは4ピクセルで効率的でなければなりません。

さらに、垂直方向に隣接するピクセルを同じマルチプロセッサで処理したいとします。その理由は、隣接する行が同じ入力データを共有するためです。ピクセル（x = 0、y = 0）が1つのMPによって処理され、ピクセル（x = 0、y = 1）が異なるMPによって処理される場合、両方のMPはそれぞれ3つのグローバルメモリ要求を発行する必要があります。両方が同じMPによって処理され、結果が適切に（暗黙的または明示的に）キャッシュされている場合は、合計で4つしか必要ありません。これは、各スレッドが複数の垂直ピクセルで動作するようにするか、blockDim.y> 1を使用することによって実行できます。

もっと一般的には、各32スレッドワープに、MP（16-48 KB、または少なくとも128x128ブロック）で使用可能なメモリをロードし、その中のすべてのピクセルその窓。

2.0より前のコンピューティング互換性のあるデバイスでは、共有メモリを使用することをお勧めします。コンピューティング互換2.0および2.1のデバイスでは、キャッシュ機能が大幅に改善されているため、グローバルメモリが問題なく動作する可能性があります。

各ワープが、入力ピクセルの各水平行に2つのキャッシュラインしかアクセスしないようにすることで、3倍ではなく、4ピクセル/ 32スレッドワープごとに。

バッファタイプとしてfloatを使用する理由はありません。メモリ帯域幅の4倍で終わるだけでなく、コードが信頼できなくなりバグが発生しやすくなります。 （浮動小数点数と整数を比較しているので、たとえばif(neighbors == 3)が正しく動作していることを確認してください）。unsigned charを使用します。より良いのは、uint8_tを使用し、unsigned charが定義されていない場合はtypedefを使用します。

最後に、実験の価値を過小評価しないでください。非常に多くの場合、CUDAコードのパフォーマンスをロジックで簡単に説明することはできません。パラメータを調整して何が起こるかを知る必要があります。

Answer 2

は、このスレッドを見て、私たちは...あそこ改善のアロット

http://forums.nvidia.com/index.php?showtopic=152757&st=60

Answer 3

TLをした; DRは：参照：http://golly.sourceforge.net

問題は、ほとんどのCUDA実装が続くということです隣人の手動カウントの脳死の考え。これは非常に遅いため、スマートなシリアルCPUの実装はそれを上回るでしょう。

GoL計算を行う唯一の賢い方法は、ルックアップテーブルを使用することです。
現在、CPU上の最も高速な実装では、将来の2x2セルを内部に取得するために、正方形の4x4 = 16ビットブロックを参照します。

この設定では、細胞がそのように同じようにレイアウトされています

0xxxxxxxx //byte0 
1xxxxxxxx //byte1 
2 etc 
3 
4 
5 
6 
7 

いくつかのビットシフトはワードに収まるように4×4ブロックを取得するために使用されると、その単語は、ルックアップテーブルを使用して検索されます。ルックアップテーブルにも単語が含まれています。このように4つの異なるバージョンの結果をルックアップテーブルに保存できるため、入力や出力に必要なビットシフト量を最小限に抑えることができます。あなたが唯一の4つの隣接スラブを見ているよう

はまた異なる世代が代わりに9 ので、同様の、互い違いに配置されていますに比べて1000倍+スピードアップで

AAAAAAAA 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
AAAAAAAA BBBBBBBB 
      BBBBBBBB 
//odd generations (A) are 1 pixel above and to the right of B, 
//even generations (B) are 1 pixels below and to the left of A. 

これだけでは結果を愚かな実装をカウントする。

そして、静的であるか、2

の周期性を持っている。そしてHashLifeがあるが、それは全く別の獣だスラブを計算していないの最適化があります。
HashLifeは、通常の実装で可能なO（n）時間ではなく、O（log n）時間でLifeパターンを生成できます。 これにより、6,366,548,773,467,669,985,195,496,000（6兆分の1）の世代を単なる秒単位で計算できます。
残念ながら、Hashlifeには再帰が必要なため、CUDAでは困難です。