問題タブ [gpu-warp]

2D配列の値を1スペース右にシフトする（行の境界をラップする）カーネルを一般化すると、ワープ同期の問題が発生します。完全なコードが添付され、以下に含まれています。

このコードは、任意の配列幅、配列の高さ、スレッドブロックの数、およびブロックあたりのスレッドの数に対して機能することを目的としています。33のスレッドサイズ（つまり、フルワープより1つ多いスレッド）を選択すると、同期されない33番目のスレッド__syncthreads()が呼び出されます。これにより、出力データに問題が発生します。この問題は、複数のワープがあり、配列の幅がスレッドの数よりも大きい場合にのみ発生します（たとえば、width = 35および34スレッドの場合）。

以下は、何が起こるかを縮小した例です（実際には、カーネルがエラーを生成するために、配列にはより多くの要素が必要になります）。

初期配列：

期待される結果：

カーネルプロデュース：

最初の行は正しく実行され（複数ある場合はブロックごとに）、後続のすべての行で最後から2番目の値が繰り返されます。私はこれを2つの異なるカード（8600GTとGTX280）でテストし、同じ結果を得ました。これが私のカーネルの単なるバグなのか、それともコードを調整しても修正できない問題なのか知りたいのですが。

完全なソースファイルは以下に含まれています。

ありがとうございました。

縦糸は32本。マルチプロセッサで 32 のスレッドが並列に実行されますか? 32 個のスレッドが並行して実行されていない場合、ワープに競合状態はありません。いくつかの例を見た後、私はこの疑問を抱きました。

ブロックに次元があるグリッドでカーネルを開始する場合：

グリッドブロックはどのようにワープに分割されますか？そのようなブロックの最初の2行は1つの縦糸を形成しますか、それとも最初の2列を形成しますか、それともこれは任意の順序ですか？

GPUコンピューティング機能を2.0と想定します。

注: この質問は、nVIDIA Compute Capability 2.1 デバイスに固有のものです。次の情報は、CUDA プログラミング ガイド v4.1 から取得されます。

コンピューティング機能 2.1 デバイスでは、各SMに整数および浮動小数点演算用の48 個のSP (コア) があります。各ワープは 32 の連続したスレッドで構成されます。各 SM には 2 つのワープ スケジューラがあります。命令が発行されるたびに、1 つのワープ スケジューラがスレッドの準備が整ったワープを選択し、コアのワープに対して2 つの命令を発行します。

私の疑問：

• 1 つのコアで 1 つのスレッドが実行されます。1 つのクロック サイクルまたは 1 つのマルチサイクル操作で、デバイスが 1 つのスレッドに 2 つの命令を発行するにはどうすればよいですか?
• これは、2 つの命令が互いに独立している必要があることを意味しますか?
• おそらくコアで異なる実行ユニットを使用するため、2つの命令をコアで並行して実行できるのでしょうか? これはまた、2 つの命令の実行が終了した後にのみワープが準備できることを意味しますか、それともそのうちの 1 つの後ですか?

次のコードは、配列内のすべての要素を各要素グループ32の最初の要素に合計します。32

__syncthreads()すべての操作が同じワープで行われるため、コード内のすべてを削除できると思いました。しかし、それらを削除すると、ガベージの結果が返されます。パフォーマンスにはあまり影響しませんが、なぜ__syncthreads()ここが必要なのか知りたいです。

私はGeForceGTX460SEを持っているので、6 SM x48CUDAコア=288CUDAコアです。1つのワープには32のスレッドが含まれ、1つのブロックで同時に（一度に）実行できるのは1つのワープのみであることが知られています。つまり、単一のマルチプロセッサ（SM）では、48コアが使用可能であっても、1つのブロック、1つのワープ、および32のスレッドのみを同時に実行できますか？

さらに、具体的なスレッドとブロックを配布する例は、threadIdx.xとblockIdx.xを使用できます。それらを割り当てるには、カーネル<<<ブロック、スレッド>>>（）を使用します。しかし、特定の数のWarpを割り当てて配布する方法と、それが不可能な場合は、なぜWarpについてわざわざ知る必要があるのでしょうか。

ワープがカーネルを終了しても、同じブロックの別のワープがまだ実行中の場合、終了したワープは、同じブロックの他のワープが終了するまでブロックされますか、または終了したワープは、別のブロックですぐに再利用できるようになります。現在のブロックのワープがまだ実行中ですか?

私の理解では、ワープは実行時にタスク スケジューラを介して定義されるスレッドのグループです。CUDA のパフォーマンスに重要な部分の 1 つは、ワープ内のスレッドの分岐です。ハードウェアがワープを構築する方法を適切に推測する方法はありますか?スレッドブロック内？

たとえば、スレッド ブロック内の 1024 スレッドでカーネルを開始しました。ワープはどのように配置されているのでしょうか。スレッド インデックスからそれを知ることができますか (または、少なくとも適切な推測を行うことができますか)。

これを行うことにより、特定のワープ内のスレッドの発散を最小限に抑えることができるためです。

オンラインレースチェックのドキュメントでは、重大度レベルに次のハザード レベルの説明があります
。

スレッドはグループで処理されるため、このステートメントは混乱を招きます。(SM はワープ全体でコードを実行します。) それらがグループで処理されない場合、それらはどのように処理されますか?
「ワープレベルプログラミング」とは？(非ワープレベルのプログラミングとは?)

GPU のワープについて質問があります。

次の構成を使用しました。

• ゲフォース210
• Cuda 機能のメジャー/マイナー: 1.2
• 2 マルチプロセッサ、8 CUDA コア/MP : 16 CUDA コア
• ワープサイズ : 32

以下は実行時間です（私はnsightを使用しました）：

Warp (=32 スレッド) が同時に実行され、2 つの MP があります。というわけで、この GPU の最大能力は 64 スレッドだと思っていましたが、16*32 スレッドはほぼ同じ時間で実行されます。ワープ スケジューラを考えると、この結果は理解できません。

私の質問は次のとおりです。

1. 16*32 スレッドが 32 スレッドとほぼ同じ時間実行されるのはなぜですか?
2. 64*32 実行時間が 32*32 の 2 倍にならない理由
3. グローバルメモリアクセスはレジスタ並みに速いと聞きました。それは正しいですか？(3.5 GPU または高価な GPU を含む)