7

特に、CUDA C での負荷分散のベスト プラクティスに関する一般的なアドバイスと説明を期待しています。

  • ワープの 1 つのスレッドが他の 31 スレッドよりも時間がかかる場合、他の 31 スレッドの完了を遅らせますか?
  • もしそうなら、予備の処理能力は別のワープに割り当てられますか?
  • なぜワープとブロックの概念が必要なのですか? ワープは 32 個のスレッドの小さなブロックにすぎないように思えます。
  • では、一般的に、カーネルへの特定の呼び出しに対して、ロード バランスは何が必要なのでしょうか?
    • 各縦糸の糸？
    • 各ブロックのスレッド？
    • すべてのブロックにまたがるスレッド?

最後に、例を挙げて、次の機能にどのような負荷分散手法を使用するかを示します。

  1. ポイントのベクトルx0がありNます：[1, 2, 3, ..., N]
  2. ポイントの5％とlogそれら（またはいくつかの複雑な関数）をランダムに選択します
  3. 結果のベクトルx1(例: [1, log(2), 3, 4, 5, ..., N]) をメモリに書き込みます。
  4. 上記の 2 つの操作を繰り返してx1生成しx2(例[1, log(log(2)), 3, 4, log(5), ..., N])、さらに 8 回繰り返して生成しx3ます ...x10
  5. 私は返すx10

どうもありがとう。

4

5 に答える 5

7

スレッドは、スケジュールが異なる 3 つのレベルにグループ化されます。Warps は SIMD を利用して計算密度を高めます。スレッド ブロックは、マルチスレッド化を利用してレイテンシの許容度を高めます。グリッドは、SM 間でロード バランシングを行うための独立した大まかな作業単位を提供します。

たて糸

ハードウェアはワープの 32 スレッドを一緒に実行します。1 つの命令を異なるデータで 32 回実行できます。スレッドが異なる制御フローを取る場合、すべてが同じ命令を実行するわけではないため、命令の実行中は 32 の実行リソースの一部がアイドル状態になります。これは、CUDA リファレンスでは制御発散と呼ばれます。

カーネルが多くの制御の相違を示す場合、このレベルで作業を再分配する価値があるかもしれません。これにより、ワープ内ですべての実行リソースをビジー状態に保つことで、作業のバランスが取れます。以下に示すように、スレッド間で作業を再割り当てできます。

// Identify which data should be processed
if (should_do_work(threadIdx.x)) {
  int tmp_index = atomicAdd(&tmp_counter, 1); 
  tmp[tmp_index] = threadIdx.x;
}
__syncthreads();

// Assign that work to the first threads in the block
if (threadIdx.x < tmp_counter) {
  int thread_index = tmp[threadIdx.x];
  do_work(thread_index); // Thread threadIdx.x does work on behalf of thread tmp[threadIdx.x]
}

ブロック内ワープ

SM では、ハードウェア スケジュールが実行ユニットにワープします。一部の命令は完了するまでに時間がかかるため、スケジューラは複数のワープの実行をインターリーブして、実行ユニットをビジー状態に保ちます。一部のワープの実行準備ができていない場合、それらはスキップされ、パフォーマンスが低下することはありません。

通常、このレベルで負荷分散を行う必要はありません。スレッド ブロックごとに十分な数のワープが利用可能であることを確認するだけで、スケジューラーはいつでも実行可能なワープを見つけることができます。

グリッド内のブロック

ランタイム システムはブロックを SM にスケジュールします。SM では複数のブロックを同時に実行できます。

通常、このレベルで負荷分散を行う必要はありません。すべての SM を数回埋めるのに十分なスレッド ブロックが使用可能であることを確認してください。一部の SM がアイドル状態で、スレッド ブロックを実行する準備ができていない場合、カーネルの最後で負荷の不均衡を最小限に抑えるために、スレッド ブロックをオーバープロビジョニングすると便利です。

于 2013-01-02T22:17:05.770 に答える
5

他の人がすでに言ったように、ワープ内のスレッドは、Single Instruction, Multiple Data (SIMD) と呼ばれるスキームを使用します。 core' は基本的に単なる浮動小数点 ALU であり、CPU コアと同じ意味での完全なコアではありません。正確な CUDA コアと命令デコーダの比率は CUDA Compute Capability のバージョンによって異なりますが、すべてこのスキームを使用しています。それらはすべて同じ命令デコーダを使用するため、スレッドのワープ内の各スレッドは、すべてのクロック サイクルでまったく同じ命令を実行します。現在実行中のコード パスをたどらないワープ内のスレッドに割り当てられたコアは、そのクロック サイクルでは何もしません。これを回避する方法はありません。これは、意図的な物理的なハードウェアの制限であるためです。したがって、1 つのワープに 32 のスレッドがあり、それらの 32 のスレッドのそれぞれが異なるコード パスをたどる場合、そのワープ内での並列処理によるスピードアップはまったくありません。これらの 32 のコード パスのそれぞれを順番に実行します。複数のスレッドが同じコード パスをたどっている場合にのみ、ワープ内の並列処理が可能であるため、ワープ内のすべてのスレッドができるだけ同じコード パスをたどるのが理想的である理由はここにあります。

ハードウェアがこのように設計されている理由は、チップ スペースを節約するためです。各コアには独自の命令デコーダがないため、コア自体が占有するチップ スペースが少なくて済みます (そして消費電力も少なくなります)。コアあたりの消費電力が少ない小さなコアを持つということは、より多くのコアをチップに搭載できることを意味します。このような小さなコアを持つことで、GPU は 1 チップあたり数百または数千のコアを持つことができますが、CPU は 4 または 8 しか持たず、同様のチップ サイズと消費電力 (および熱放散) レベルを維持します。SIMD とのトレードオフは、より多くの ALU をチップに詰め込み、より多くの並列処理を実現できることですが、これらの ALU がすべて同じコード パスを実行している場合にのみ高速化が得られます。このトレードオフが GPU で非常に高度に行われる理由 s は、3D グラフィック処理に含まれる計算の多くが、単純に浮動小数点行列の乗算であることです。結果の行列の各出力値を計算するプロセスは、データが異なるだけで同じであるため、SIMD は行列の乗算に適しています。さらに、各出力値は他のすべての出力値とは完全に独立して計算できるため、スレッドが相互に通信する必要はまったくありません。ちなみに、同様のパターン (および行列の乗算自体でさえも) は、科学および工学のアプリケーションでもよく見られます。これが、GPU の汎用処理 (GPGPU) が生まれた理由です。

于 2013-01-03T15:30:00.730 に答える
4

ワープの 1 つのスレッドが他の 31 スレッドよりも時間がかかる場合、他の 31 スレッドの完了を遅らせますか?

はい。ワープで分岐が発生するとすぐに、スケジューラはすべての分岐分岐を取得して、1 つずつ処理する必要があります。現在実行されているブランチにないスレッドの計算能力は失われます。CUDA Programming Guide を確認できます。正確に何が起こるかについて非常によく説明されています。

もしそうなら、予備の処理能力は別のワープに割り当てられますか?

いいえ、残念ながらそれは完全に失われています。

なぜワープとブロックの概念が必要なのですか? ワープは 32 個のスレッドの小さなブロックにすぎないように思えます。

ワープは最適なパフォーマンスを実現するために SIMD (単一命令、複数データ) でなければならないため、ブロック内のワープは完全に分岐する可能性がありますが、他のリソースを共有します。(共有メモリ、レジスタなど)

では、一般的に、カーネルへの特定の呼び出しに対して、ロード バランスは何が必要なのでしょうか?

ロードバランスという言葉は適切ではないと思います。常に十分な数のスレッドが常に実行されていることを確認し、ワープ内での分岐を回避してください。繰り返しますが、CUDA プログラミング ガイドは、そのようなことを読むのに適しています。

次に例を示します。

m=0..N*0.05 で m スレッドを実行し、それぞれが乱数を選択し、「複雑な関数」の結果を x1[m] に入れることができます。ただし、大規模な領域でグローバル メモリからランダムに読み取ることは、GPU で実行できる最も効率的な方法ではないため、完全にランダムにする必要があるかどうかについても検討する必要があります。

于 2013-01-02T21:49:50.263 に答える
2

他の人は、理論的な質問に対して良い答えを提供しています.

あなたの例では、次のように問題を再構築することを検討してください。

  1. 点のベクトルxを持つN：[1, 2, 3, ..., N]
  2. のすべての要素について複雑な関数を計算しx、 を生成しyます。
  3. を介しyて生成するのサブセットをランダムにサンプリングします。y0y10

ステップ 2 は、その値が必要かどうかを考慮せずに、すべての入力要素を 1 回だけ操作します。ステップ 3 のサンプリングが置換なしで行われる場合、これは、実際に必要な要素数の 2 倍を計算することになりますが、制御の発散なしですべてを計算し、すべてのメモリ アクセスが一貫性を持つことを意味します。これらは多くの場合、計算そのものよりも GPU の速度を左右する重要な要素ですが、これは複雑な関数が実際に何を行っているかによって異なります。

ステップ 3 は非コヒーレントなメモリ アクセス パターンになるため、GPU で実行する方がよいのか、それとも CPU に戻してそこでサンプリングを行う方が速いのかを判断する必要があります。

次の計算が何であるかに応じて、ステップ 3 を再構成して、各要素の [0,N) の整数をランダムに描画することができます。値が [N/2,N) にある場合、次の計算では無視します。[0,N/2) にある場合は、その値をその仮想y*配列 (または計算に適したもの) のアキュムレータに関連付けます。

于 2013-01-03T01:51:29.190 に答える
1

あなたの例は、削減を示す本当に良い方法です。

I have a vector x0 of N points: [1, 2, 3, ..., N]
I randomly pick 50% of the points and log them (or some complicated function) (1)
I write the resulting vector x1 to memory (2)
I repeat the above 2 operations on x1 to yield x2, and then do a further 8 iterations to  yield x3 ... x10 (3)
I return x10 (4)

|x0| と言います。= 1024 で、ポイントの 50% を選択します。

最初の段階は、グローバル メモリから読み取る必要がある唯一の段階である可能性があります。その理由を説明します。

512 個のスレッドがメモリから 512 個の値を読み取り (1)、共有メモリに格納します (2)。次に、ステップ (3) で 256 個のスレッドが共有メモリからランダムな値を読み取り、共有メモリにも格納します。スレッドが 1 つになるまでこれを行い、グローバル メモリに書き戻します (4)。

これをさらに拡張するには、最初のステップで 256 のスレッドで 2 つの値を読み取るか、128 のスレッドで 4 つの値を読み取るなど...

于 2013-01-03T22:53:22.683 に答える