第7章: GPU设计中的流式多处理器设计
流式多处理器（SM）是NVIDIA GPU体系结构的基础构建块。每个SM包含一组CUDA核心,这些 CUDA 核心以SIMT（单指令,多线程）方式执行指令。 SM负责管理和调度warp,处理分支发散,并提供对共享内存和缓存的快速访问。在本章中,我们将探讨 SM 的微体系结构,包括其管线、warp 调度机制、寄存器文件设计以及共享内存和 L1 缓存组织。
SM微体系结构和管线
SM是一个高度并行和流水线处理器,旨在高效地并发执行数百个线程。图7.1展示了NVIDIA Volta体系结构中一个SM的简化框图。
指令缓存
|
v
Warp调度器
|
v
调度单元(4个warp)
| | | |
v v v v
CUDA核心(FP64/FP32/INT)
CUDA核心(FP64/FP32/INT)
CUDA核心(FP64/FP32/INT)
...
Tensor核心
Tensor核心
...
载入/存储单元
载入/存储单元
...
特殊功能单元
^
|
寄存器文件(64 KB)
^
```共有メモリ / L1キャッシュ (96 KB)
図7.1: NVIDIA VoltaアーキテクチャにおけるSMの簡略化されたブロック図
SMの主要なコンポーネントは以下の通りです:
-
命令キャッシュ: 頻繁にアクセスされる命令を格納し、レイテンシの低減とスループットの向上を図ります。
-
ワープスケジューラ: 実行準備が整ったワープを選択し、利用可能な実行ユニットにディスパッチします。
-
ディスパッチユニット: 1サイクルに最大4つのワープの命令をフェッチ、デコードし、適切な実行ユニットにディスパッチします。
-
CUDA コア: プログラマブルな実行ユニットで、幅広い整数演算と浮動小数点演算をサポートします。Voltaのeach SMには64個のCUDAコアが含まれます。
-
テンサーコア: ディープラーニングやAIワークロードの高速化を目的とした専用の実行ユニットです。Voltaのeach SMには8個のテンサーコアが含まれます。
-
ロード/ストアユニット: グローバルメモリ、共有メモリ、キャッシュへのメモリ操作(ロードとストア)を処理します。
-
特殊関数ユニット: 超越関数やその他の複雑な数学演算を実行します。
-
レジスタファイル: スレッドプライベートなレジスタへの高速アクセスを提供します。Voltaのeach SMには64 KBのレジスタファイルがあります。
-
共有メモリ / L1キャッシュ: ソフトウェア管理キャッシュ(共有メモリ)またはハードウェア管理L1データキャッシュとして使用できる設定可能なメモリ空間です。
SMパイプラインは、複数のワープを同時に実行し、メモリレイテンシを隠蔽することで、スループットを最大化するように設計されています。図7.2は、SMパイプラインの簡略化された図です。
命令フェッチ
|
v
命令デコード
|
v
オペランド収集
|
v
実行 (CUDAコア、テンサーコア、ロード/ストアユニット、特殊関数ユニット)
|
v
書き戻し図7.2: 簡略化されたSMパイプライン
パイプラインの各段階は以下の通りです:
-
命令フェッチ: ワープスケジューラは実行準備が整ったワープを選択し、ここは日本語の翻訳です。コードの部分は翻訳されていません。
-
命令フェッチ: 現在実行中のワープのための次の命令がキャッシュから取り出されます。
-
命令デコード: フェッチされた命令がデコードされ、命令の種類、オペランド、および宛先レジスタが特定されます。
-
オペランド収集: 命令の実行に必要なオペランドがレジスタファイルやシared メモリから収集されます。
-
実行: 命令がCUDAコア、テンソルコア、ロード/ストアユニット、または特殊関数ユニットなどの適切な実行ユニットで実行されます。
-
書き戻し: 実行の結果がレジスタファイルやシared メモリに書き戻されます。
高いパフォーマンスを達成するために、SMは以下の手法を使ってリソースの利用を最大化し、待機時間を隠蔽しています:
- デュアル発行: SMは1サイクルあたり2つの独立した命令をワープに発行できるため、命令レベルの並列性が高まります。
- パイプラインの実行ユニット: 実行ユニットはパイプラインされているため、前の操作が完了する前に新しい操作を開始できます。
- 待機時間の隠蔽: SMはサイクルごとにワープを切り替えることができ、メモリアクセスや長い待ち時間の命令の待機時間を他のワープの命令を実行することで隠蔽できます。
例7.1は、2つのベクトルの要素単位の加算を行うシンプルなCUDAカーネルを示しています。
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
c[tid] = a[tid] + b[tid];
}
}例7.1: ベクトル加算のCUDAカーネル。
この例では、カーネルの各スレッドが入力ベクトターaとbの対応する要素の和を計算し、その結果を出力ベクトターcに格納します。SMはこのカーネルを実行するためにCUDAコアに各スレッドを割り当て、使用可能なコアでワープのスレッドをスケジューリングします。ロード/ストアユニットはグローバルメモリから入力データをフェッチし、結果を書き戻すために使用されます。
ワープのスケジューリングと分岐の処理
Efこちらが日本語の翻訳です。コードの部分は翻訳していません。
高効率なワープスケジューリングは、SM (ストリーミングマルチプロセッサ) のパフォーマンスを最大化するために重要です。ワープスケジューラは、実行準備ができたワープを選択し、利用可能な実行ユニットにディスパッチする責任があります。ワープスケジューラの主な目的は、常にワープが実行可能な状態にあることで、実行ユニットを忙しく保つことです。
SMは2つのレベルのワープスケジューリングメカニズムを採用しています:
-
ワープスケジューリング: ワープスケジューラは、ラウンドロビンやoldestrFirst などのスケジューリングポリシーに基づいて、実行準備ができたワープを選択します。選択されたワープは、利用可能な実行ユニットにディスパッチされます。
-
命令スケジューリング: 各ワープ内で、SMは命令の依存関係と実行ユニットの利用可能性に基づいて命令をスケジューリングします。SMは、命令レベルの並列性を最大化するために、同一ワープから複数の独立した命令を同時に発行できます。
図7.3は、2つのレベルのワープスケジューリングメカニズムを示しています。
Warp Pool
Warp 1 (Ready)
Warp 2 (Waiting)
Warp 3 (Ready)
...
Warp N (Ready)
|
v
Warp Scheduler
|
v
Dispatch Unit
|
v
Execution Units図7.3: 2つのレベルのワープスケジューリングメカニズム
ワープスケジューリングの主な課題の1つは、分岐の発散への対処です。SIMTモデルでは、同一ワープ内のすべてのスレッドが同期して同じ命令を実行します。しかし、ワープが分岐命令(if-else文など)に遭遇すると、一部のスレッドがif-pathを、他のスレッドがelse-pathを取ることがあります。これは分岐の発散と呼ばれる状況です。
分岐の発散に対処するために、SMはプリケーション (predication) と呼ばれる技術を使用しています。ワープが発散する分岐に遭遇すると、SMはその両方のパスを順次実行し、それぞれのパスに属さないスレッドをマスクします。その後、プリケートレジスタを使用して結果を結合し、各スレッドが正しい結果を得られるようにします。
例 7.2は、分岐の発散が見られるCUDAカーネルを示しています。以下は、提供されたマークダウンファイル「branch」の日本語訳です。コードについては、コメントのみを翻訳しています。
__global__ void divergentKernel(int *data, int *result) {
// スレッドIDを取得
int tid = blockIdx.x * blockDim.x + threadIdx.x;
// データが0より大きい場合
if (data[tid] > 0) {
// 結果をデータの2倍にする
result[tid] = data[tid] * 2;
} else {
// 結果をデータの3倍にする
result[tid] = data[tid] * 3;
}
}例7.2: 分岐が発散するCUDAカーネル
この例では、分岐条件 data[tid] > 0 によって、ワープ内の一部のスレッドがif-pathを、他のスレッドがelse-pathを実行することになります。SMはこの発散を処理するために、両方のパスを順次実行し、各パスで非アクティブなスレッドをマスクします。
図7.4は、発散したスレッドを持つワープの述語化プロセスを示しています。
ワープ (32スレッド)
スレッド1: data[1] = 5, result[1] = 10
スレッド2: data[2] = -3, result[2] = -9
...
スレッド32: data[32] = 7, result[32] = 14
発散した分岐:
if (data[tid] > 0) {
result[tid] = data[tid] * 2;
} else {
result[tid] = data[tid] * 3;
}
述語化:
ステップ1: マスクを使ってif-pathを実行
スレッド1: result[1] = 10
スレッド2: (マスクされている)
...
スレッド32: result[32] = 14
ステップ2: マスクを使ってelse-pathを実行
スレッド1: (マスクされている)
スレッド2: result[2] = -9
...
スレッド32: (マスクされている)
最終結果:
スレッド1: result[1] = 10
スレッド2: result[2] = -9
...
スレッド32: result[32] = 14図7.4: 発散したスレッドを持つワープの述語化プロセス
述語化を使うことで、SMは明示的な分岐命令や制御フロー発散なしにブランチ発散を処理できます。しかし、発散したブランチはまだパフォーマンスに影響を及ぼし、SMが両方のパスを順次実行しなければならないため、効果的な並列性が低下します。
レジスタファイルとオペランドコレクタ
レジスタファイルは、高速なスレッドプライベートレジスタへのアクセスを提供する、SMの重要なコンポーネントです。各SMには多数のアクティブなスレッドをサポートし、ワープ間の効率的なコンテキスト切り替えを可能にする大規模なレジスタファイルがあります。 NVIDIAのVolta アーキテクチャにおいて、各SMには64 KBのレジスタファイルが組織化されており、32個の2 KBのバンクから構成されています。レジスタファイルは、大量の並行スレッドをサポートするため、高帯域幅と低待時間のアクセスを提供するよう設計されています。
バンクコンフリクトを最小限に抑え、パフォーマンスを改善するため、SMはオペランド収集と呼ばれる手法を採用しています。オペランド収集ユニットは、レジスタファイルのバンクからオペランドを収集し、実行ユニットに供給します。オペランド収集を使うことで、SMはバンクコンフリクトの影響を軽減し、実行ユニットの利用効率を高めることができます。
図7.5は、SMのレジスタファイルとオペランド収集ユニットの簡略図を示しています。
レジスタファイル (64 KB)
バンク 1 (2 KB)
バンク 2 (2 KB)
...
バンク 32 (2 KB)
|
v
オペランド収集ユニット
|
v
実行ユニット図7.5: SMのレジスタファイルとオペランド収集ユニット
オペランド収集ユニットは、複数のインストラクションおよび複数のワープからオペランドを収集することで、SMが1サイクルで異なるワープからのインストラクションを実行ユニットに発行できるようにします。これにより、レジスタファイルアクセスの待機時間を隠蔽し、SMの全体的なスループットを向上させることができます。
例7.3は、2つのベクトルの内積を計算するCUDAカーネルを示しています。
__global__ void dotProduct(float *a, float *b, float *result, int n) {
// 部分和を保持する共有メモリ
__shared__ float partialSum[256];
// スレッドのインデックス
int tid = threadIdx.x;
// 処理するデータのインデックス
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 部分和を0で初期化
partialSum[tid] = 0;
// 割り当てられたデータを処理
while (i < n) {
partialSum[tid] += a[i] * b[i];
i += blockDim.x * gridDim.x;
}
// 同期
__syncthreads();
// 部分和を集約
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
partialSum[tid] += partialSum[tid + s];
}
__syncthreads();
}
// 結果を書き込み
if (tid == 0) {
result[blockIdx.x] = partialSum[0];
}
}この例では、各スレッドが割り当てられた部分のデータを処理し、部分和を計算します。入力ベクトルの要素から成る. 部分和は共有メモリ配列 partialSum に格納される. すべてのスレッドが自分の部分和を計算し終えたら, 並列の縮小演算が行われ, 最終的な内積結果が得られる.
オペランドコレクターは, この例において重要な役割を果たす. 共有メモリアクセスと算術演算のためのオペランドを効率的に収集することで, バンク競合を避け, 演算ユニットの活用率を高めることができる.
結論
ストリーミングマルチプロセッサ (SM) は, 現代のGPUアーキテクチャにおける中核的な計算ユニットである. その設計は, 細かい粒度のマルチスレッディング, SIMTの実行, そしてオペランド収集の効率化を組み合わせることで, スループットの最大化とメモリ待機時間の隠蔽に重点を置いている.
SMの主要なコンポーネントには, ワープのスケジューリングを行うワープスケジューラ, ブランチの分岐と収束を処理するSIMTスタック, スレッドプライベートレジスタへの高速アクセスを提供するレジスタファイルとオペランドコレクター, そして低待機時間のデータ共有とリユースを可能にする共有メモリとL1キャッシュがある.
GPUアーキテクチャの進化とともに, ブランチ分岐の処理, ワープスケジューリング, レジスタファイルの設計などの分野での研究が, 将来のGPUの性能と効率の向上に不可欠となるだろう. ダイナミックワープ形成, スレッドブロック圧縮, オペランドリユースキャッシュなどの新しい手法は, SMの機能を大幅に強化し, 並列コンピューティングワークロードにおける新たな性能レベルを実現する可能性を秘めている.