0

「1 ブロック x 32 スレッド」の構成で開始したいカーネルがあります。並列性を高めるために、「1 ブロック x 32 スレッド」よりも大きな「作業パッケージ」を実行する代わりに、複数のストリームを開始したいと考えています。データがネットワークから来るプログラムで GPU を使用したいと考えています。より大きな「作業パッケージ」が利用可能になるまで待ちたくありません。コードは次のようになります。

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
  - copy data GPU -> Host [cudaMemcpyAsync(.., stream i)]
}

実際のコードはもっと複雑ですが、シンプルに保ちたいと思います (15 個の CPU スレッドが GPU を使用します)。

コードは機能しますが、ストリームは期待どおりに同時に実行されません。GTX 480 には 15 の SM があり、各 SM には 32 のシェーダー プロセッサがあります。カーネルを 15 回起動すると、15 個のストリームすべてが並行して実行されると思いますが、そうではありません。私は Nvidia Visual Profiler を使用しており、並行して実行される最大 5 つのストリームがあります。多くの場合、1 つのストリームのみが実行されます。パフォーマンスは本当に悪いです。

「64 ブロック x 1024 スレッド」構成で最良の結果が得られます。代わりに「32 ブロック x 1024 スレッド」構成を使用すると、2 つのストリームが次々に実行され、パフォーマンスが低下します。Cuda Toolkit 5.5 と Ubuntu 12.04 を使用しています。

誰かがこれが事実である理由を説明し、背景情報を教えてもらえますか? 新しい GPU ではうまく動作するはずですか? データをバッファリングしたくないタイムクリティカルなアプリケーションで GPU を使用する最良の方法は何ですか? おそらくこれは不可能ですが、解決策に近づけるテクニックを探しています。

ニュース：

さらに調査を行いました。問題は、最後の cudaMemcpyAsync(..) (GPU->ホスト コピー) 呼び出しです。削除すると、すべてのストリームが同時に実行されます。問題は、スライド 21 のhttp://on-demand.gputechconf.com/gtc-express/2011/presentations/StreamsAndConcurrencyWebinar.pdfに示されていると思います。Fermiには 2 つのコピー キューがあると言われていますが、これはテスラカードとクワドロカードですよね？問題は、GTX 480 にはコピー キューが 1 つしかなく、すべてのコピー コマンド (ホスト -> GPU および GPU -> ホスト) がこの 1 つのキューに入れられることだと思います。すべてがノンブロッキングで、最初のスレッドの GPU->host memcopy は、他のスレッドの host->GPU memcopy 呼び出しをブロックします。ここでいくつかの観察：

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
}

-> 動作: ストリームは同時に実行されます

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
  - sleep(10)
  - copy data GPU -> Host [cudaMemcpyAsync(.., stream i)]
}

-> 動作: ストリームは同時に実行されます

Thread(i=0..14) {
  - copy data Host -> GPU [cudaMemcpyAsync(.., stream i)]
  - run kernel(stream i)
  - cudaStreamSynchronize(stream i)
  - copy data GPU -> Host [cudaMemcpyAsync(.., stream i)]
}

- >動作しません!!! cudaStreamSynchronize が copy-queue に入れられているのではないでしょうか?

誰かがこの問題の解決策を知っていますか。ブロッキング カーネル コールのようなものがクールでしょう。カーネルが終了したら、最後の cudaMemcpyAsync() (GPU->device) を呼び出す必要があります。

Edit2: ここに私の問題を明確にする例があります:単純にするために、2つのストリームがあります:

Stream1:
------------
HostToGPU1
kernel1
GPUToHost1

Stream2:
------------
HostToGPU2
kernel2
GPUToHost2

最初のストリームが開始されます。HostToGPU1 が実行され、kernel1 が起動され、GPUToHost1 が呼び出されます。kernel1 が実行されているため、GPUToHost1 はブロックします。その間に Stream2 が開始されます。HostToGPU2 が呼び出され、Cuda はそれをキューに入れますが、カーネル 1 が終了するまで GPUToHost1 がブロックされるため、実行できません。現在、データ転送はありません。Cuda は GPUToHost1 を待つだけです。したがって、私の考えは、kernel1 が終了したときに GPUToHost1 を呼び出すことでした。これが、カーネルが終了したときに GPUToHost1 が呼び出されるため、sleep(..) で動作する理由であると思われます。CPU スレッドを自動的にブロックするカーネル起動はクールです。GPUToHost1 はキューでブロックされていません (その時点で他にデータ転送がない場合ですが、私の場合はデータ転送に時間がかかりません)。

4

1 に答える 1

2

カーネルの同時実行は、Linux で最も簡単に確認できます。

良い例と簡単なテストについては、並行カーネルのサンプルを参照してください。

カーネル間の良好な並行性には、一般にいくつかのことが必要です。

  • 並行カーネルをサポートするデバイスなので、cc 2.0 以降のデバイス
  • 複数のカーネルが実際に実行できるように、ブロック数およびその他のリソース使用量 (レジスター、共有メモリー) の点で十分に小さいカーネル。より大きなリソース要件を持つカーネルは、通常、連続して実行されていることが観察されます。これは予期される動作です。
  • 同時実行を可能にするためのストリームの適切な使用

さらに、並行カーネルは多くの場合、コピー/計算のオーバーラップを意味します。コピー/計算のオーバーラップを機能させるには、次のことを行う必要があります。

  • 十分なコピー エンジンを備えた GPU を使用している。一部の GPU には 1 つのエンジンがあり、いくつかの GPU には 2 つのエンジンがあります。GPU にエンジンが 1 つある場合、1 つのコピー操作 (つまり、一方向) をカーネル実行とオーバーラップさせることができます。コピー エンジンが 2 つある場合 (GeForce GPU に 1 つある場合)、コピーの両方向をカーネル実行とオーバーラップさせることができます。
  • GPU グローバル メモリとの間でコピーされるデータには固定 (ホスト) メモリを使用します。これは、オーバーラップするコピー操作のターゲット (との間) になります。
  • ストリームを適切に使用し、関連する API 呼び出しの必要な非同期バージョン (例:cudaMemcpyAsync

小さい 32x1024 カーネルが同時に実行されないという観察に関しては、これはおそらくリソースの問題 (ブロック、レジスタ、共有メモリ) であり、多くのオーバーラップが妨げられています。最初のカーネルに GPU 実行リソースを占有するのに十分なブロックがある場合、最初のカーネルが終了するかほとんど終了するまで、追加のカーネルが実行を開始することを期待するのは賢明ではありません。

編集:質問の編集と以下の追加コメントへの対応。

はい、GTX480にはコピー「キュー」が1つしかありません（これは私の回答で明示的に言及しましたが、コピー「エンジン」と呼びました）。任意の時点で cudaMemcpy... 操作を1 つしか実行できないため、実際にデータを移動できるのは 1 つの方向 (H2D または D2H) だけであり、 1 つのcudaMemcpy...しか表示されません。操作は、任意のカーネルとオーバーラップします。そして、そのストリームに対して以前に発行されたすべてcudaStreamSynchronizeの CUDA 操作が完了するまで、ストリームを待機させます。

cudaStreamSynchronize最後の例にある は必要ないことに注意してください。ストリームには 2 つの実行特性があります。

  1. 同じストリームに対して発行された cuda 操作 (API 呼び出し、カーネル呼び出し、すべて)は、 API の使用やその他の考慮事項 に関係なく、常に順次実行されます。Async
  2. 必要なすべての要件が満たされていると仮定すると、個別のストリームに発行された cuda 操作は、互いに非同期で実行されます。

項目 1 により、最後のケースでは、最後の「copy Data GPU->Host」操作は、呼び出しがなくても、そのストリームに対して発行された前のカーネル呼び出しが完了するまで開始されませんcudaStreamSynchronize。したがって、その呼び出しを取り除くことができると思います。つまり、リストした2番目のケースは最後のケースと変わらないはずであり、2番目のケースではスリープ操作も必要ありません。同じストリームに発行された cudaMemcpy... は、そのストリーム内の以前のすべての cuda アクティビティが終了するまで開始されません。これがストリームの特徴です。

EDIT2:ここで進歩が見られるかどうかはわかりません。こちらの GTC preso で指摘された問題(スライド 21) は有効な問題ですが、追加の同期操作を挿入しても回避できません。また、「ブロッキング カーネル」がそれを助けることも、機能ではないこともあります。 1 つまたは 2 つのコピー エンジンを使用する場合。別のストリームでカーネルを発行するが、他の介在する cuda 操作なしで順番に発行する場合、その危険性が存在します。これに対する解決策は、次のスライドで指摘されているように、カーネルを連続して発行しないことです。これは、2 番目のケースとほぼ同じです。これをもう一度述べます：

  • ケース2が良好な同時実行性を与えることを確認しました
  • その場合のスリープ操作は、データの整合性のために必要ありません

問題を示す短いサンプル コードを提供したい場合は、おそらく他の発見を行うことができます。

于 2013-08-21T01:49:24.700 に答える