5

＃blocksとblockSizeの選び方については多くの議論がありましたが、それでも何かが足りません。私の懸念の多くはこの質問に対処しています：CUDAブロック/ワープ/スレッドはどのようにCUDAコアにマッピングされますか？ （説明を簡単にするために、十分なperThreadおよびperBlockメモリがあります。メモリ制限はここでは問題になりません。）

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal);

1）SMをできるだけビジーに保つためnThreadsに、の倍数に設定する必要がありwarpSizeます。本当ですか？

2）SMは、一度に1つのカーネルしか実行できません。つまり、そのSMのすべてのHWcoreはkernelAのみを実行しています。（kernelAを実行しているHWcoreもあれば、kernelBを実行しているHWcoreもありません。）したがって、実行するスレッドが1つしかない場合は、他のHWcoreを「無駄にしている」ことになります。本当ですか？

3）ワープスケジューラーの発行がwarpSize（32スレッド）単位で機能し、各SMに32個のHWコアがある場合、SMは十分に活用されます。SMに48個のHWcoreがあるとどうなりますか？スケジューラーが32のチャンクで作業を発行しているときに、48のコアすべてを完全に利用し続けるにはどうすればよいですか？（前の段落が当てはまる場合、スケジューラーがHWcoreサイズの単位で作業を発行した方がよいのではないでしょうか？）

4）ワープスケジューラが一度に2つのタスクをキューに入れているように見えます。そのため、現在実行中のカーネルがストールまたはブロックしたときに、2番目のカーネルがスワップインされます（明確ではありませんが、ここのキューは2カーネル以上の深さだと思います）。これは正しいですか？

5）ハードウェアの上限がブロックあたり512スレッド（nThreadsMax）である場合、512スレッドのカーネルが1つのブロックで最速で実行されることを意味するわけではありません。（繰り返しになりますが、問題ではありません。）512スレッドのカーネルを1つだけでなく多くのブロックに分散させると、パフォーマンスが向上する可能性が高くなります。ブロックは1つまたは複数のSMで実行されます。本当ですか？

5a）小さい方がいいと思いますが、どれだけ小さくしてもnBlocks大丈夫ですか？問題は、その価値をどのように選択するかというnBlocksことです。（必ずしも最適ではありません。）を選択するための数学的アプローチはありますかnBlocks、それとも単に試行錯誤です。

4

2 に答える 2

5

1）はい。

2）CC 2.0-3.0デバイスは、最大16グリッドを同時に実行できます。各SMは8ブロックに制限されているため、完全な同時実行性を実現するには、デバイスに少なくとも2つのSMが必要です。

3）はい、ワープスケジューラは一度にワープを選択して発行します。CUDAコアの概念を忘れてください。それらは無関係です。レイテンシーを隠すには、高い命令レベルの並列性または高い占有率が必要です。CC1.xの場合は>25％、CC> = 2.0の場合は>50％にすることをお勧めします。一般に、CC 3.0は、スケジューラーが2倍になるため、2.0デバイスよりも高い占有率を必要としますが、SMあたりのワープは33％しか増加しません。Nsight VSE Issue Efficiency実験は、命令とメモリのレイテンシを隠すのに十分なワープがあるかどうかを判断するための最良の方法です。残念ながら、ビジュアルプロファイラーにはこのメトリックがありません。

4）ワープスケジューラアルゴリズムは文書化されていません。ただし、スレッドブロックがどのグリッドから発生したかは考慮されません。CC 2.xおよび3.0デバイスの場合、CUDA作業ディストリビューターは、次のグリッドからブロックを配布する前に、グリッドからすべてのブロックを配布します。ただし、これはプログラミングモデルによって保証されていません。

5）SMをビジー状態に保つには、デバイスを満たすのに十分なブロックが必要です。その後、適度な占有率に達するのに十分なワープがあることを確認する必要があります。大きなスレッドブロックを使用することには、長所と短所の両方があります。一般に、大きなスレッドブロックは、使用する命令キャッシュが少なく、キャッシュ上のフットプリントが小さくなります。ただし、大きなスレッドブロックは同期スレッドで停止し（選択するワープが少なくなるため、SMの効率が低下する可能性があります）、同様の実行ユニットで命令を実行し続ける傾向があります。開始するには、スレッドブロックごとに128または256スレッドを試すことをお勧めします。スレッドブロックが大きい場合と小さい場合の両方に十分な理由があります。5a）占有計算機を使用します。スレッドブロックサイズが大きすぎると、レジスタによって制限されることがよくあります。

于 2012-06-07T20:10:17.420 に答える
3

あなたの質問に一つずつ答えてみようと思います。

  1. それは正しいです。
  2. 「HWcores」とはどういう意味ですか？ステートメントの最初の部分は正しいです。

  3. NVIDIA Fermi Compute Architectureホワイトペーパーによると：「SMはワープと呼ばれる32の並列スレッドのグループでスレッドをスケジュールします。各SMは2つのワープスケジューラと2つの命令ディスパッチユニットを備えており、2つのワープを同時に発行および実行できます。Fermiのデュアルワープスケジューラは2つのワープ、および各ワープから16コアのグループ、16のロード/ストアユニット、または4つのSFUに1つの命令を発行します。ワープは独立して実行されるため、Fermiのスケジューラーは命令ストリーム内から依存関係をチェックする必要はありません。

    さらに、NVIDIA Kepplerアーキテクチャホワイトペーパーには、「Keplerのクアッドワープスケジューラは4つのワープを選択し、ワープごとに2つの独立した命令を各サイクルでディスパッチできます」と記載されています。

    したがって、「過剰な」コアは、一度に複数のワープをスケジュールすることによって使用されます。

  4. ワープスケジューラは、異なるカーネルではなく、同じカーネルのワープをスケジュールします。

  5. 正確ではありません。共有メモリが存在する場所であるため、各ブロックは単一のSMにロックインされます。
  6. これはトリッキーな問題であり、カーネルの実装方法によって異なります。より重要な問題のいくつかを説明しているVasilyVolkovによるnVidiaウェビナーの低占有率でのパフォーマンスの向上をご覧ください。ただし、主に、 CUDA Occupancy Calculatorを使用して、占有率を向上させるためにスレッド数を選択することをお勧めします。
于 2012-06-07T15:29:58.780 に答える