1

私はCUDAの初心者です。

これまでのところ、各SMには（スレッドの）8つのブロックがあることを学びました。配列内の要素に2を掛けるという単純な仕事があるとしましょう。ただし、データはスレッドよりも少なくなっています。

スレッドの「テール」を切り取ってアイドル状態にすることができたので、問題ありません。しかし、私が正しく理解していれば、これは、一部のSMが100％の作業を取得し、一部の部分（またはまったく機能しない）を取得することを意味します。

そこで、どのSMが特定のスレッドで実行されているかを計算し、各SMが同じ量の作業を行うように計算したいと思います。

そもそもそれが理にかなっていることを願っています:-)もしそうなら、どのSMが与えられたスレッドを実行しているかを計算する方法は？または-現在のSMのインデックスとそれらの総数？言い換えれば、SM用語でthreadDim/threadIdxと同等です。

アップデート

コメントとしては長すぎました。

ロバート、答えてくれてありがとう。私はすべてを消化しようとしていますが、これが私が行うことです-私は「大きな」配列を持っており、値を乗算し*2て出力配列に格納する必要があります（ウォームアップとして;ところで、私が行うすべての計算は数学的に正しいです）。したがって、最初にこれを1ブロック、1スレッドで実行します。罰金。次に、各乗算が1つのスレッドで1回だけ行われるように作業を分割しようとしました。その結果、私のプログラムの実行速度は約6倍遅くなります。GPUに関する情報を取得し、使用するブロックとスレッドの数を計算し、単一の乗算の代わりに各スレッド内で、配列内のオフセットを計算するためだけに約10の追加の乗算があるため、理由がわかります。スレッド。一方では、その望ましくない動作を変更する方法を見つけようとします。他方では、スレッドの「テール」をSM間で均等に広げたいと思います。

言い換えれば、私は間違っているかもしれませんが、これを解決したいと思います。私には1G*2の小さな仕事があります（それだけです）-1Kスレッドで1Kブロックを作成するか、1スレッドで1Mブロックを作成するか、1Mスレッドで1ブロックを作成する必要があります。これまでのところ、GPUプロパティを読み取り、分割、分割し、グリッド/ブロックの各次元の最大値（または計算するデータがない場合は必要な値）を盲目的に使用します。

コード

sizeは入力および出力配列のサイズです。一般に：

output_array[i] = input_array[i]*2;

必要なブロック/スレッドの数を計算します。

size_t total_threads = props.maxThreadsPerMultiProcessor
                       * props.multiProcessorCount;
if (size<total_threads)
    total_threads = size;

size_t total_blocks = 1+(total_threads-1)/props.maxThreadsPerBlock;

size_t threads_per_block = 1+(total_threads-1)/total_blocks;

props.maxGridSize持っているとprops.maxThreadsDim私は同様の方法でブロックとスレッドの寸法を計算します-fromtotal_blocksとthreads_per_block。

そして、キラー部分は、スレッド（スレッドの「内部」）のオフセットを計算します。

size_t offset = threadIdx.z;
size_t dim = blockDim.x;
offset += threadIdx.y*dim;
dim *= blockDim.y;
offset += threadIdx.z*dim;
dim *= blockDim.z;
offset += blockIdx.x*dim;
dim *= gridDim.x;
offset += blockIdx.y*dim;
dim *= gridDim.y;

size_t chunk = 1+(size-1)/dim;

これで、現在のスレッドの開始オフセットと、乗算用の配列（チャンク）内のデータ量がわかりました。grimDim.zAFAIKは常に1なので、上記は使用しませんでしたね。

4

1 に答える 1

6

やろうとするのは珍しいことです。あなたがCUDAの初心者であることを考えると、そのような質問は、問題を不適切に解決しようとしていることを示しているように思われます。あなたが解決しようとしている問題は何ですか？SMXとSMYで特定のスレッドを実行している場合、問題はどのように役立ちますか？マシンから最大のパフォーマンスが必要な場合は、すべてのスレッドプロセッサとSMがアクティブになり、実際にはすべてに「十分な作業」があるように作業を構成します。GPUは、レイテンシを隠すためにオーバーサブスクライブされたリソースに依存しています。

CUDA初心者として、あなたの目標は次のとおりです。

  • ブロックとスレッドの両方で十分な作業を作成する
  • メモリに効率的にアクセスします（これは主に合体と関係があります-あなたはそれを読むことができます）

「各SMに同じ量の作業がある」ことを確認することにメリットはありません。グリッドに十分な数のブロックを作成すると、各SMの作業量はほぼ等しくなります。これはスケジューラーの仕事です、あなたはスケジューラーにそれをさせるべきです。十分な数のブロックを作成しない場合、最初の目的は、実行する作業をさらに作成または見つけることであり、利益をもたらさないブロックごとの派手な作業分解図を考え出すことではありません。

たとえば、Fermi GPUの各SMには、32個のスレッドプロセッサがあります。メモリアクセスなどによる不可避のマシンストールが存在する場合でもこれらのプロセッサをビジー状態に保つために、マシンは、ストールが発生したときに別のスレッドのワープ（32）にスワップすることでレイテンシを隠し、処理を続行できるように設計されています。 。これを容易にするために、SMごとに多数の利用可能なワープを用意するようにしてください。これは、次のようにすることで容易になります。

  • グリッド内の多数のスレッドブロック（GPU内のSMの数の少なくとも6倍）
  • スレッドブロックごとに複数のワープ（おそらく少なくとも4〜8ワープ、つまりブロックごとに128〜256スレッド）

（Fermi）SMは常に一度に32スレッドを実行しているため、GPU内のSMの数の32倍より少ないスレッドがある場合、マシンは十分に活用されていません。私の問題全体が、たとえば20スレッドのみで構成されている場合、GPUを利用するように設計されていないだけであり、それらの20スレッドを複数のSM/スレッドブロックに分割してもそれほどメリットはありません。

編集：あなたはあなたのコードを投稿したくないので、私はさらにいくつかの提案やコメントをします。

  1. いくつかのコードを変更しようとしましたが、実行速度が遅いことがわかり、間違った結論にジャンプしました（私は思います）。
  2. おそらく、 vectoraddのような単純なコード例に精通している必要があります。各要素を乗算しているわけではありませんが、構造は近いです。単一のスレッドを使用してこのベクトル追加を実行すると、実際に高速に実行される方法はありません。この例を研究すると、配列要素に2を掛けて実行するように拡張する簡単な方法が見つかると思います。
  3. あなたが概説したように、誰もブロックごとのスレッドを計算しません。まず、ブロックあたりのスレッド数は32の倍数である必要があります。次に、ブロックあたりのスレッド数を開始点として選択し、そこから他の起動パラメーターを作成するのが通例です。その逆ではありません。大きな問題の場合は、ブロックあたり256または512スレッドから始めて、その計算を省略します。
  4. 選択したスレッドブロックサイズに基づいて、他の起動パラメータ（グリッドサイズ）を作成します。問題は本質的に1Dであるため、1Dスレッドブロックの1Dグリッドが出発点として適しています。この計算がx次元の最大ブロックに関してマシンの制限を超える場合は、各スレッドループで複数の要素を処理するか、（1Dスレッドブロックの）2Dグリッドに拡張することができます。
  5. オフセットの計算は不必要に複雑です。配列を処理するための比較的単純なオフセット計算でスレッドのグリッドを作成する方法については、ベクトル追加の例を参照してください。
于 2013-02-10T22:21:44.703 に答える