私のCUDAコードは、4つの配列で動作する必要があります（平均/標準に縮小し、ヒストグラムを計算します）。各2048浮動小数点数は長く、以前のカーネルからデバイスメモリにすでに保存されています。
一般に、マルチプロセッサと少なくとも同じ数のブロックを起動することをお勧めします。ただし、この場合、これらの各配列を 1 つのブロックの共有メモリにロードできるため、起動できるブロックは 4 つだけです。
これは「GPU をビジー状態に保つ」には程遠いですが、より多くのブロックを使用する場合、グローバル メモリを介してより多くのブロック間通信を行う必要があります。グローバルメモリの内外。
このような状況で並列化する最良の方法は何ですか?
この例は、Thrustを使用した1回の削減ですべての「要約統計量」を計算する方法を示しています。別の例は、を使用してヒストグラムを計算する方法を示していますthrust::sort。
ブロック間通信は CUDA では推奨されません。また、Fermi はカーネルの同時実行をサポートするため、将来的に占有率が高くなることはそれほど重要ではなくなります。したがって、パフォーマンスが許容できないほど低い場合を除き、今のところ占有率を低くしておくことをお勧めします。
一度に 2048 個の float をすべて必要とする必要はないと思います。削減する必要がある場合は、配列をさまざまな部分に分割し、ブロック実行の最後に結果をマージできます。サンプルコードを表示できますか?
実行する作業の量は比較的少ないため、おそらく 4 つのブロックに固執する必要があります。前/後続のカーネルのデータを GPU に対してローカルに保つという利点はまだあります。
Fermiは並行カーネルを許可し、このカーネルの実行中に次のカーネルを開始して残りの SM を占有できるため、まさにこのケースが最もメリットがあります。ただし、これは 2 つのカーネル間に依存関係がないことを前提としています。当然、前のカーネルが終了する前に、前のカーネルの結果に依存する新しいカーネルを開始することはできません。