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

スレッドブロックは、逐次または並列に実行できるスレッドのグループを表すプログラミング抽象概念です。プロセスとデータのマッピングを改善するために、スレッドはスレッドブロックにグループ化されます。スレッドブロック内のスレッド数は、以前はアーキテクチャによってブロックあたり合計512スレッドに制限されていましたが、2010年3月現在、Compute Capability 2.x以降では、ブロックに最大1024スレッドを含めることができます。同じスレッドブロック内のスレッドは、同じストリームマルチプロセッサ上で実行されます。^{[ 1 ]}同じブロック内のスレッドは、共有メモリ、バリア同期、またはアトミック操作などの他の同期プリミティブを 介して相互に通信できます。

複数のブロックが結合されてグリッドを形成します。同じグリッド内のすべてのブロックには、同じ数のスレッドが含まれます。ブロック内のスレッド数には制限がありますが、グリッドは、多数のスレッドブロックを並列処理し、利用可能なすべてのマルチプロセッサを使用する必要がある計算に使用できます。

CUDAは、高水準言語が並列処理を活用するために使用できる並列コンピューティングプラットフォームおよびプログラミングモデルです。CUDAでは、カーネルはスレッドの助けを借りて実行されます。スレッドはカーネルの実行を表す抽象的なエンティティです。カーネルは、特定のデバイス上で実行するためにコンパイルされる関数です。マルチスレッドアプリケーションは、同時に実行される多数のスレッドを使用して並列計算を構成します。各スレッドにはインデックスがあり、メモリアドレスの位置を計算したり、制御の決定を下したりするために使用されます。

寸法

CUDAは、ホストデバイス上のアプリケーションプログラムを実行するための異種プログラミングモデルに基づいて動作します。CUDAの実行モデルはOpenCLに似ています。このモデルでは、アプリケーションの実行はホストデバイス（通常はCPUコア）上で開始されます。ホストデバイスはスループット重視のデバイス、つまり並列計算を実行するGPUコアです。これらの並列実行にはカーネル関数が使用されます。カーネル関数が実行されると、制御はホストデバイスに戻され、シリアル実行が再開されます。

多くの並列アプリケーションは多次元データを扱うため、スレッド ブロックを 1D、2D、または 3D のスレッド配列に編成すると便利です。グリッド内のブロック間では通信や連携ができないため、グリッド内のブロックは独立して実行できる必要があります。「カーネルの起動時に、スレッド ブロックあたりのスレッド数とスレッド ブロックの数が指定され、これによって起動される CUDA スレッドの総数が定義されます。 」 ^{[ 2 ]}ブロックの x、y、z の最大次元は 1024、1024、64 であり、x × y × z ≤ 1024 (ブロックあたりの最大スレッド数) となるように割り当てる必要があります。^{[ 3 ]ブロックは、x、y、z 次元でそれぞれ最大 2 31} -1、65,535、65,535 ブロックの 1 次元、2 次元、または 3 次元グリッドに編成できます。^{[ 3 ]}ブロックあたりの最大スレッド数とは異なり、最大グリッド寸法とは異なるグリッドあたりのブロック数の制限はありません。

インデックス作成

1Dインデックス

CUDA のすべてのスレッドは特定のインデックスに関連付けられているため、配列内のメモリ位置を計算してアクセスできます。

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

CUDAプログラミングにおけるスレッド階層^{[ 4 ]}

スレッドインデックス 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 ] ; } }

2Dインデックス

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

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

ハードウェアの観点

スレッドの階層構造について述べましたが、スレッド、スレッドブロック、グリッドは基本的にプログラマーの視点からの説明であることに留意してください。スレッドブロックの本質を完全に理解するには、ハードウェアの観点から理解することが重要です。ハードウェアは、同じ命令を実行するスレッドをワープにグループ化します。複数のワープが1つのスレッドブロックを構成します。複数のスレッドブロックはストリーミングマルチプロセッサ（SM）に割り当てられます。複数のSMがGPUユニット全体を構成します（このユニットがカーネルグリッド全体を実行します）。

GPUのスレッドブロックに対するプログラマーの視点とハードウェアの視点の相関図^{[ 7 ]}

ストリーミングマルチプロセッサ

GPUの各アーキテクチャ（KeplerやFermiなど）は、複数のSM（ストリーミング・マルチプロセッサ）で構成されています。これらは、低クロックレートをターゲットとし、キャッシュ容量の小さい汎用プロセッサです。SMは複数のスレッドブロックを並列に実行できます。あるスレッドブロックの実行が完了すると、次のスレッドブロックを順次実行します。一般的に、SMは命令レベルの並列処理をサポートしますが、分岐予測はサポートしません。^{[ 8 ]}

ストリーミングマルチプロセッサとそのリソースの図解^{[ 9 ]}

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

• 実行コア。(単精度浮動小数点ユニット、倍精度浮動小数点ユニット、特殊関数ユニット (SFU))。
• キャッシュ:

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

• ワープのスケジューラ。(特定のスケジューリング ポリシーに基づいてワープに指示を発行します)。
• 相当数のレジスタ。(SM は一度に多数のアクティブ スレッドを実行する可能性があるため、数千のレジスタが必要になります。)

ハードウェアはスレッドブロックをSMにスケジュールします。通常、SMは複数のスレッドブロックを同時に処理できます。SMは最大8個のスレッドブロックを保持できます。スレッドIDは、それぞれのSMによってスレッドに割り当てられます。

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

NvidiaのFermiマイクロアーキテクチャに実装されたダブルワープスケジューラの図解^{[ 10 ]}

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

ワープ

ハードウェア側では、スレッドブロックは「ワープ」で構成されています。（この用語はウィービング^{[ 12 ]}に由来します。 ）ワープとは、スレッドブロック内の32個のスレッドの集合です。従来、これらのスレッドは「ロックステップ」で実行されること（ワープ内のすべてのスレッドが同時に命令を実行すること）が保証されており、さらに重要な点として、すべてのメモリ位置へのアクセスは、ワープスレッドのすべて、あるいは全く実行されないことが保証されていました。この動作は、簡単にデッドロックにつながる可能性がありました（例えば、ループ内でif分岐を使用する場合など）。しかし、Voltaアーキテクチャ以降、より細粒度のロックを介してワープ内のデータ交換が可能になりました。^{[ 13 ] [ 14 ]}これらのスレッドは、SMによって順次選択されます。^{[ 15 ]}

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

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

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

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

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

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

参考文献

1. 「第 4 章 ハードウェア実装、スレッド ブロックのスレッドは 1 つのマルチプロセッサ上で同時に実行され、複数のスレッド ブロックは 1 つのマルチプロセッサ上で同時に実行できます」。
2. 「CUDAスレッドモデル」www.olcf.ornl.gov . 2016年9月23日時点のオリジナルよりアーカイブ。 2016年9月21日閲覧。
3. 「CUDAツールキットドキュメント：機能と技術仕様」 . docs.nvidia.com . 2022年5月24日閲覧。
4. 「CUDAプログラミングにおけるスレッド階層」 。 2016年9月21日閲覧。
5. カーク、デイビッド、ヒュー、ウェンメイ W (2010年1月28日). 『超並列プロセッサのプログラミング：実践的アプローチ』 .
6. 「スレッドインデックス作成チートシート」（PDF） . 2016年9月21日閲覧。
7. 「スレッドの最適化（メイランド大学）」（PDF）。
8. Wilt, Nicholas (2013). CUDAハンドブック：GPUプログラミングの包括的なガイド.
9. 「スレッドの最適化（メイランド大学）」（PDF）。
10. 「スレッドの最適化（メイランド大学）」（PDF）。
11. 「CUDA による GPU コンピューティング 講義 2 - CUDA メモリ」(PDF)。
12. 「Parallel Thread Execution ISA Version 6.0」。開発者ゾーン：CUDAツールキットドキュメント。NVIDIA Corporation。2017年9月22日。2017年10月28日時点のオリジナルよりアーカイブ。 2017年10月27日閲覧。
13. 「1. Voltaチューニングガイド — Voltaチューニングガイド13.0ドキュメント」 . docs.nvidia.com . 2025年8月5日閲覧。
14. Nvidia. 「Cuda C++プログラミングモデルV13」(PDF) . p. 142. 2025年8月24日閲覧。
15. 「CUDAワープレベルプリミティブの使用」。Nvidia 。 2018年1月15日。 2020年4月8日閲覧。NVIDIA GPUは、ワープと呼ばれるスレッドのグループをSIMT（単一命令複数スレッド）方式で実行します。
16. 「CUDA のメモリ問題と CUDA の実行スケジュール」(PDF)。
17. 「命令フェッチとメモリスケジューリングの GPU パフォーマンスへの影響」(PDF)。
18. 「CAWS: GPGPU ワークロードのクリティカリティを考慮したワープ スケジューリング」(PDF)。