スレッドブロック（CUDAプログラミング）

CUDAでは、各スレッドは特定のインデックスに関連付けられており、それによって配列内のメモリ位置を計算したりアクセスしたりすることができます。

512 個の要素を持つ配列の例を考えてみましょう。組織構造の 1 つは、512 個のスレッドを持つ単一のブロックを持つグリッドを採用することです。512 個の要素を持つ 2 つの配列 A と B の要素ごとの乗算によって構成される 512 個の要素を持つ配列 C があるとします。各スレッドにはインデックス i があり、A と B の i 番目の要素の乗算を実行し^、その結果を C の i^番目の要素に格納します。i は、ブロックが 1 つしかないため、この場合は 0 である blockIdx、ブロックが 512 個の要素を持つため、この場合は 512、および各ブロックに対して 0 から 511 まで変化する threadIdx を使用して計算されます。

スレッドインデックス i は、次の式で計算されます 。

${\displaystyle i=blockIdx.x*blockDim.x+threadIdx.x}$

blockIdx.x は x 次元ブロック識別子です

blockDim.x はブロックの寸法の x 次元です

threadIdx.x はスレッド識別子の x 次元です

したがって、「i」は0から511までの値を取り、配列全体を網羅します。

1024 より大きい配列の計算を考慮する場合、それぞれ 1024 スレッドを持つ複数のブロックを使用できます。2048 個の配列要素を持つ例を考えてみましょう。この場合、それぞれ 1024 スレッドを持つ 2 つのスレッド ブロックがあります。したがって、スレッド識別子の値は 0 から 1023 まで変化し、ブロック識別子は 0 から 1 まで変化し、ブロックの次元は 1024 になります。したがって、最初のブロックは 0 から 1023 までのインデックス値を取得し、最後のブロックは 1024 から 2047 までのインデックス値を取得します。

したがって、各スレッドはまずアクセスする必要のあるメモリのインデックスを計算し、次に計算を進めます。スレッドを使用して配列 A と B の要素を並列に追加し、結果を配列 C に格納する例を考えてみましょう。スレッド内の対応するコードは次のとおりです 。^{[ 5 ]}

__global__ void vecAddKernel ( float * A , float * B , float * C , int n ) { int index = blockIdx . x * blockDim . x + threadIdx . x ; if ( index < n ) { C [ index ] = A [ index ] + B [ index ] ; } }

同様に、特に複雑なグリッドでは、ブロックIDとスレッドIDは、グリッドの形状に応じて各スレッドで計算する必要があります。2次元ブロックを持つ2次元グリッドを考えてみましょう。スレッドIDとブロックIDは、次の式で計算されます 。

${\displaystyle blockId=blockIdx.x+blockIdx.y*gridDim.x;}$${\displaystyle threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;}$^{[ 6 ]}

GPU の各アーキテクチャ (例えばKeplerやFermi ) は、複数の SM (ストリーミング マルチプロセッサ) で構成されています。これらは、クロック レートが低く、キャッシュが小さい汎用プロセッサです。SM は、複数のスレッド ブロックを並列に実行できます。スレッド ブロックの 1 つが実行を完了するとすぐに、次のスレッド ブロックを直列に取ります。一般的に、SM は命令レベルの並列処理をサポートしますが、分岐予測はサポートしません。^{[ 8 ]}

この目的を達成するために、SMには以下が含まれます。^{[ 8 ]}

• 実行コア。（単精度浮動小数点演算ユニット、倍精度浮動小数点演算ユニット、特殊機能ユニット（SFU））
• キャッシュ:

1. L1キャッシュ（メモリアクセス遅延を低減するため）。
2. 共有メモリ（スレッド間でデータを共有するため）。
3. 定数キャッシュ（読み取り専用メモリからの読み取りをブロードキャストするため）。
4. テクスチャキャッシュ（テクスチャメモリからの帯域幅を集約するため）。

• ワープ用スケジューラ。（これらは、特定のスケジューリングポリシーに基づいてワープに指示を発行するためのものです。）
• 相当数のレジスタ。（SMは同時に多数のアクティブなスレッドを実行する可能性があるため、数千個のレジスタが必須です。）

ハードウェアはスレッドブロックをSM（ストリーミングマルチプロセッサ）に割り当てます。一般的に、1つのSMは複数のスレッドブロックを同時に処理できます。1つのSMには最大8つのスレッドブロックを含めることができます。スレッドIDは、それぞれのSMによってスレッドに割り当てられます。

SMがスレッドブロックを実行する際、そのスレッドブロック内のすべてのスレッドが同時に実行されます。したがって、SM内のスレッドブロックのメモリを解放するには、ブロック内のすべてのスレッドの実行が完了していることが不可欠です。各スレッドブロックは、ワープと呼ばれるスケジューリング単位に分割されます。ワープについては、次のセクションで詳しく説明します。

SMのワープスケジューラは、命令の発行時にどのワープを優先するかを決定します。^{[ 11 ]}ワープの優先順位付けポリシーの一部については、次のセクションでも説明されています。

ハードウェア側では、スレッドブロックは「ワープ」で構成されます。（この用語は織物から来ています。^{[ 12 ]} ）ワープは、スレッドブロック内の32個のスレッドのセットです。従来、これらのスレッドは「同期して」（ワープ内のすべてのスレッドが同時に命令を実行する）実行され、重要なことに、すべてのメモリ位置にすべてのワープスレッドまたはどのワープスレッドもアクセスしないことが保証されていました。この動作は、（ループ内でif分岐を使用することによって）デッドロックに容易につながる可能性がありました。しかし、Voltaアーキテクチャ以降、より細かい粒度のロックによるワープ内データ交換が可能になりました。^{[ 13 ] [ 14 ]}これらのスレッドはSMによって直列に選択されます。^{[ 15 ]}

マルチプロセッサ（SM）上でスレッドブロックが起動されると、そのすべてのワープは実行が終了するまで常駐します。したがって、新しいブロックのすべてのワープに必要な数の空きレジスタと、新しいブロックに必要な十分な空き共有メモリが確保されるまで、新しいブロックはSM上で起動されません。

32 個のスレッドからなるワープが命令を実行している場合を考えてみましょう。オペランドの 1 つまたは両方が準備できていない場合 (たとえば、グローバル メモリからまだフェッチされていない場合)、コンテキスト スイッチングと呼ばれるプロセスが実行され、制御が別のワープに渡されます。^{[ 16 ]}特定のワープから切り替えるとき、そのワープのすべてのデータはレジスタ ファイル内に残るため、オペランドが準備できたときにすぐに再開できます。命令に未解決のデータ依存関係がない場合、つまり両方のオペランドが準備できている場合、対応するワープは実行準備完了とみなされます。実行可能なワープが複数ある場合、親 SM はワープスケジューリング ポリシーを使用して、次にフェッチされる命令をどのワープに渡すかを決定します。

実行可能なワープをスケジュールするためのさまざまなポリシーについては、以下で説明します。^{[ 17 ]}

1. ラウンドロビン（RR） - 命令はラウンドロビン方式でフェッチされます。RRは、SM（ストリーミングマルチプロセッサ）が常に稼働状態を維持し、メモリ遅延によるクロックサイクルの無駄をなくします。
2. 最も最近フェッチされていない命令(LRF) - このポリシーでは、最も長い間フェッチされていない命令のワープが、命令のフェッチにおいて優先されます。
3. フェア（FAIR）^{[ 17 ]} - このポリシーでは、スケジューラはすべてのワープに対してフェッチされる命令の数に関して「公平な」機会が与えられるようにします。フェッチされた命令の数が最小のワープに命令をフェッチします。
4. スレッドブロックベースのCAWS ^{[ 18 ]} (クリティカル性を考慮したワープスケジューリング) - このスケジューリングポリシーの重点は、スレッドブロックの実行時間を改善することにあります。実行に最も時間がかかるワープに多くの時間リソースを割り当てます。最も重要なワープに優先順位を付けることで、このポリシーはスレッドブロックの終了を早め、リソースをより早く利用可能にします。

従来のCPUスレッドコンテキストの切り替えでは、割り当てられたレジスタ値とプログラムカウンタをオフチップメモリ​​（またはキャッシュ）に保存および復元する必要があるため、ワープコンテキストの切り替えよりもはるかに負荷の高い処理となります。ワープのレジスタ値（プログラムカウンタを含む）はすべてレジスタファイルに保持され、スレッドブロック内のすべてのワープ間で共有される共有メモリ（およびキャッシュ）もそのまま保持されます。

ワープアーキテクチャの利点を最大限に活用するには、プログラミング言語と開発者は、メモリアクセスを統合する方法と、制御フローの分岐を管理する方法を理解する必要があります。ワープ内の各スレッドが異なる実行パスを辿ったり、各スレッドが著しく異なるメモリにアクセスしたりすると、ワープアーキテクチャの利点が失われ、パフォーマンスが著しく低下します。