人工知能システムの基盤となる計算処理において、GPU（Graphics Processing Unit）は不可欠な役割を担っている。本稿では、GPUのハードウェア構造とNVIDIA GPUアーキテクチャの観点から、その動作原理を詳細に解説する。NVIDIA GPUはFermiからBlackwellアーキテクチャに至るまで長年にわたる発展を遂げており、特にTensor CoreやNVLinkなどのAI関連技術は目覚ましい進化を遂げている。
本稿では、まずCPUとGPUのアーキテクチャの違いを明らかにし、続いてAX+Y演算を題材にGPUによる並列処理の具体的な仕組みを考察する。さらに、GPUにおける重要な設計思想である並列性（Parallelism）と並行性（Concurrency）の概念についても詳しく解説する。最後に、GPUのメモリアクセスとスレッド管理について深く掘り下げ、その設計哲学を理解することを目的とする。
CPUとGPUの本質的な違い
GPUとは何か、そしてCPUとGPUはどのような点で異なるのかを明確にすることが、GPUの動作原理を理解するための第一歩となる。
CPU（Central Processing Unit：中央処理装置）は、オペレーティングシステムやアプリケーションの実行に必要なあらゆる計算処理を担当する。汎用性の高い処理能力を持ち、多様なデータ型を取り扱う必要がある。論理的な判断に伴う分岐処理や割り込み処理が頻繁に発生するため、CPUの内部構造は非常に複雑多岐にわたる。
一方、GPU（Graphics Processing Unit：グラフィックス処理装置）は、並列実行時に複雑な数学的演算を効率的に処理することに特化した半導体チップである。当初はゲームやアニメーションにおけるグラフィックスレンダリング処理のために設計されたが、現在ではその用途は大きく拡張され、深層学習や科学技術計算など幅広い分野で利用されている。両者は共に、コア、メモリ、制御ユニットといった基本的な構成要素を持っている。
CPUとGPUのアーキテクチャにおける主要な違いは、以下の4点に集約される。
並列処理能力の観点から 살펴보면、CPUは少数ながら高性能な演算ユニット（ALU：Arithmetic Logic Unit）を備えており、順序性のある処理任务に適している。非常に少ないクロックサイクルで算術演算を完了することが可能であり、高いクロック周波数を維持できる。複雑な制御ロジックユニットは、プログラムに複数の分岐がある場合に分岐予測機能を提供することで、効率的な処理を実現する。パイプライン技術は複数の処理を並列に実行することでプログラム実行時間を短縮する。対照的に、GPUの制御ユニットは複数のメモリアクセスを統合する機能を持ち、膨大な数の演算ユニット（ALU）とスレッド（Thread）で構成されている。多量のALUにより高い計算スループットを実現でき、余剰に設計されたスレッドはメモリアクセスの遅延を効果的に隠蔽できるため、複数のタスクを同時に処理し、大規模な並列計算任务に最適化されている。
メモリアーキテクチャの点では、CPUのダイ面積の大部分がキャッシュメモリ（Cache）に割り当てられている。多量のキャッシュメモリは後続でアクセスされる可能性のあるデータを保持することで、アクセス遅延を低減する。これに対し、GPUのキャッシュメモリは少量でスレッドごとにサービスを提供する設計となっている。多くのスレッドが同一データへのアクセスを必要とする場合、キャッシュメモリはこれらのアクセスを統合した後、DRAM（Dynamic Random Access Memory）にアクセスし、取得データを該当するスレッドに分配する。GPUにおいては大量のレジスタが準備されており、これらが多数のスレッドを同時に実行することを可能にしている。
命令セットの観点からは、CPUの命令セットはより汎用的であり、多様なタイプのタスクを実行するのに適している。GPUの命令セットは主としてグラフィックス処理と汎用計算（GPGPU：General-Purpose computing on Graphics Processing Units）に使用される。CPUは異なる命令セット間で迅速に切り替えることができるが、GPUは大量の同一命令を取得し、高速で連続的に実行する設計となっている。
消費電力と発熱の面では、CPUは比較的低い消費電力で済み、発熱要件も相対的に低い。GPUは高度に並列化された特性から、通常消費電力が高く、安定した動作を維持するためにはより優れた散热システムが求められる。
以上の特性から、CPUはオペレーティングシステムやデータ分析など、順序実行が必要な処理任务に最適であり、GPUはグラフィックス処理や深層学習など、大規模な並列計算を必要とする処理任务に適している。ヘテロジニアスシステム（異種混合システム）では、GPUとCPUを組み合わせることで、それぞれの利点を最大限に引き出すことが一般的である。
GPUは当初、グラフィックス画像やビデオエンコード/デコード関連作業を処理するために設計された。GPUとCPUの最も本質的な違いは、GPUの設計目標がスループット（Throughput）の最大化にある点にある。個々のタスクを実行する速度よりも、同時に実行可能なタスク数（Parallelism：並列度）を重視する。一方CPUは、遅延（Latency）と並行性（Concurrency）をより重視する傾向がある。
CPUの最適化目標は、可能な限り低い遅延で可能な限り快速にタスクを完了しつつ、タスク間での素早い切り替え能力を維持することである。その本質は、シーケンシャル（順序的）なタスク処理にある。GPUの最適化はすべてスループットの向上に寄与しており、可能な限り多くのタスクを一度にGPU内部に推送し、GPUは多数のコアによってこれらのタスクを並列に処理する。
プロセッサの帯域幅（Bandwidth）、遅延（Latency）、スループット（Throughput）
帯域幅は、プロセッサが処理できる最大のデータ量または命令数を表し、Kb、Mb、Gbなどの単位で表される。遅延は、プロセッサが命令を実行したりデータを処理したりするまでに要する時間であり、1データユニットを転送するために必要な時間を意味し、ms、s、min、hなどの単位で表される。スループットは、一定時間内に、ある位置から別の位置へ移動するデータ量を表し、bps（Bits Per Second：毎秒ビット数）、Mbps（Megabits Per Second：毎秒メガビット数）、Gbps（Gigabits Per Second：毎秒ギガビット数）などの単位で表される。例えば、10秒目に20ビットのデータを転送した場合、t=10におけるスループットは20 bpsとなる。
帯域幅の問題を解決することは、遅延の問題を解決するよりも容易である。スレッド数はスループットに比例し、スループットがほぼ帯域幅に等しい状態はチャネルの使用効率が高いことを意味する。プロセッサシステムの設計における目標は、帯域幅を向上させる前提下での遅延隠蔽の実現であり、実用的なプロセッサシステムを構成することである。
並列性と並行性
並列性（Parallelism）と並行性（Concurrency）は、コンピュータサイエンスの分野で頻繁に取り上げられる2つの概念である。両者とも同時に複数のタスクを処理する能力に関わるが、その具体的な意味と適用には相違点がある。
並列性（Parallelism） は、複数のタスクや操作を同時に実行することを指し、複数の処理ユニット上で同時に行われる。コンピュータシステムにおいて、これらの処理ユニットはマルチコアプロセッサ、マルチスレッド、分散システムなどを含みうる。並列計算は大量のデータや複雑な計算を処理する必要がある場合に、システムの性能と効率を大幅に向上させることができる。例えば、コンピュータプログラムは複数のプロセッサコア同時に実行することで、全体の計算速度を向上させることができる。
並行性（Concurrency） は、システムが複数のタスクや操作を同時に処理できることを指すが、必ずしも同時に実行されるわけではない。並行システムでは、タスク間が交代で実行され、タイムスライス輪番（Time Slicing）やイベント駆動（Event-Driven）などの手法によって実現される。並行性は、システムの応答性能とリソース使用効率を向上させるために使用され、特に多数の短時間タスクを処理する必要がある場合に有効である。例えば、Webサーバは複数のクライアントリクエストを同時に処理し、並行処理によってシステムのスループットを向上させることができる。
以上から、並列性と並行性の主な違いは明確になる。並列性は複数のタスクを同時に実行することを意味し、同時性と並列処理能力を重視し、計算性能と効率の向上に使用される。並行性はシステムが複数のタスクを同時に処理できることを意味し、タスク間の交代実行とリソース共有を重視し、システムの応答性能とリソース使用効率の向上に使用される。
実際の応用においては、並列性と並行性は通常組み合わせて使用され、具体的な要件とシステムの特性に応じて適切な技術と戦略を選択する。同時に、並列性と並行性の概念を理解することは、複雑なコンピュータシステムとアプリケーションの設計と最適化に不可欠である。実際のハードウェア動作の過程では、マルチスレッドを活用したループ展開によって全体のハードウェア利用率を向上させる、これがGPUの最も本質的な原理である。
3種類のチップを例として挙げ、ハードウェアの制約下で一般的にいくつのスレッドを実行可能かを比較する。比較結果として、スレッドの必要数（Threads required）、利用可能なスレッド数（Threads available）、スレッド比率（Thread Ratio）を示し，到底いくつのスレッドがあればメモリ遅延の問題を解決できるかを検証する。表中可以看到几个关键的数据：
- GPU（NVIDIA A100）の遅延はCPU（AMD Rome 7742、Intel Xeon 8280）と比較して数倍高い
- GPUのスレッド数はCPUの20〜30倍である
- GPUの利用可能なスレッド数はCPUの100倍以上である
スレッド比率を計算した結果、GPUは5.6倍、CPUは1.2〜1.3倍となった，这也是GPU最重要的设计点，它拥有非常多的线程为大规模任务并行去设计。
| AMD Rome 7742 | Intel Xeon 8280 | NVIDIA A100 | |
|---|---|---|---|
| Memory B/W(GB/sec) | 204 | 143 | 1555 |
| DRAM Latency(ns) | 122 | 89 | 404 |
| Peak bytes per latency | 24,888 | 12,727 | 628,220 |
| Memory Efficiency | 0.064% | 0.13% | 0.0025% |
| Threads required | 1,556 | 729 | 39,264 |
| Threads available | 2,048 | 896 | 221,184 |
| Thread Ratio | 1.3X | 1.2X | 5.6X |
CPUとGPUの典型的なアーキテクチャを比較すると、GPUは巨大なスループット処理装置として機能することができる。一部のスレッドはデータの到着を待ち、一部のスレッドは計算処理への活性化を待ち、一部のスレッドは計算処理中である。GPUのハードウェア設計エンジニアは、すべてのハードウェアリソースを遅延時間の削減ではなく、より多くのスレッドを追加することに使用している。命令実行の遅延を削減するのではなく、スレッド数を増やすことで隠蔽する設計思想である。
対照的に、CPUは遅延重視の処理装置と比喻できる。その主な仕事は1つのスレッド内で全ての処理を完了することであり、遅延の問題を解決するために十分なスレッドを使用することを前提としている。因此CPUの硬件设计者会把所有资源和重心都投入到减少延迟上面，因此CPU的线程比只有一点多倍，这也是SIMD（Single Instruction, Multiple Data）とSIMT（Single Instruction, Multiple Threads）アーキテクチャ間の最大の違いである。CPUはスレッド数を増やすことで問題を解決するのではなく、反対のアプローチでスレッドの実行速率と効率を最適化する，这就是CPU跟GPU之间最大的区别，也是它们的本质区别。
SIMDとSIMT
SIMD（Single Instruction, Multiple Data）アーキテクチャは、同時に複数のデータに対して同一の操作を実行するものであり、ベクトル化演算に適している。例えば、複数の要素を含む配列に対して、SIMDアーキテクチャはすべての要素に対して同時に同じ操作を実行することで、計算効率を向上させる。SSE（Streaming SIMD Extensions）やAVX（Advanced Vector Extensions）が代表的なSIMDアーキテクチャである。
SIMT（Single Instruction, Multiple Threads）アーキテクチャは、同時に複数のスレッドを実行するものであり、各スレッドは異なる命令を実行できるが、これらのスレッドは通常同じプログラムを実行する。このアーキテクチャはGPUにおける並列計算に使用される。CUDA（Compute Unified Device Architecture）やOpenCLは、SIMTアーキテクチャをサポートするプログラミングモデルである。
SIMDはデータ並列計算に、SIMTはタスク並列計算に適している。実際の応用においては、具体的な計算要件とハードウェア環境に応じて適切なアーキテクチャを選択することで、計算性能を向上させることができる。
GPUの動作原理
基本的な動作原理
まず、AX+Yという加算演算の例を通じてGPUの動作原理を理解する。AX+Yのサンプルコードを以下に示す。
void saxpy(double alpha, double* x, double* y, int n)
{
for (int i = 0; i < n; ++i)
{
y[i] = alpha * x[i] + y[i];
}
}
このサンプルコードには2つのFLOPS演算が含まれている。乗算（Multiply）と加算（Add）であり、各計算操作でメモリから2つのデータ、x[i]とy[i]を読み出し、線形演算を実行してy[i]に格納する。この乗算と加算を融合した操作は、FMA（Fused Multiply and Add：融合積和演算）と呼ばれる。
O(n)の計算量において、nの値に応じてn回の反復計算を実行する。CPUでは、AX+Yプログラムが命令シーケンスに従ってシリアルに実行される。Intel Xeon 8280を例にとると、メモリ帯域幅は131 GB/s、メモリ遅延は89 nsであり、これは8280チップが89 nsの間に11,659バイト（byte）のデータを転送できることを意味する。AX+Yでは89 nsの間に16バイト（C/C++におけるdoubleデータ型のメモリサイズは8 bytes）のデータを転送するため、メモリ使用率はわずか0.14%（16/11,659）に過ぎず、メモリバスの99.86%の時間がアイドル状態であることがわかる。
異なるプロセッサでAX+Yを実行した場合のメモリ使用率を比較すると、AMD Rome 7742、Intel Xeon 8280、NVIDIA A100のいずれにおいても、AX+Yプログラムのメモリ使用率は非常に低く、基本的0.14%以下である。
| AMD Rome 7742 | Intel Xeon 8280 | NVIDIA A100 | |
|---|---|---|---|
| Memory B/W(GB/sec) | 204 | 131 | 1555 |
| DRAM Latency(ns) | 122 | 89 | 404 |
| Peak bytes per latency | 24,888 | 11,659 | 628,220 |
| Memory Efficiency | 0.064% | 0.14% | 0.0025% |
上述のAX+Yプログラムは並列性と計算の線形性を十分に活用していないため、並行性（Concurrency）を活用したループ展開を行う。以下に改善後のコードを示す。
void saxpy_unrolled(int n, double alpha, double* x, double* y)
{
for (int i = 0; i < n; i += 8)
{
y[i + 0] = alpha * x[i + 0] + y[i + 0];
y[i + 1] = alpha * x[i + 1] + y[i + 1];
y[i + 2] = alpha * x[i + 2] + y[i + 2];
y[i + 3] = alpha * x[i + 3] + y[i + 3];
y[i + 4] = alpha * x[i + 4] + y[i + 4];
y[i + 5] = alpha * x[i + 5] + y[i + 5];
y[i + 6] = alpha * x[i + 6] + y[i + 6];
y[i + 7] = alpha * x[i + 7] + y[i + 7];
}
}
このコードは0から7までのデータを一度に処理し、8回の反復を1回のイタレーションで実行する。1回の転送で16バイトのデータを処理するため、同じIntel Xeon 8280チップでは、89 nsの間に11,659 ÷ 16 = 729回のリクエストを実行する。このようなプログラムの改善は、並行性を活用することでバス全体を忙碌状態に保つことを可能にする。
しかしながら、実際のアプリケーションシナリオでは以下の制約が存在する。
- コンパイラがループ全体を100回以上展開することは稀である
- 1つのスレッドが実行できる命令数は限られており、非常に多くの並行数を処理することはできない
- 1つのスレッドで700以上の計算負荷を直接処理することは困難である
由此可见，虽然并发的操作能够一次性执行更多的指令流水线操作，但是同样架构也会受到限制和约束。
次に、Z = AX + Yを並列処理で展開する方法を以下に示す。
void saxpy_parallel(int n, double alpha, double* x, double* y)
{
#pragma omp parallel for
for (int i = 0; i < n; i++)
{
y[i] = alpha * x[i] + y[i];
}
}
この方法は並列処理によるループ展開であり、並列プロセッサまたは複数のスレッドを使用してAX+Y操作を実行することで、总线忙碌状態を維持し、1回あたり729回の反復を実行することが可能となる。並行処理と比較した場合の特徴は以下の通りである。
- 各スレッドが独立して関連する計算を担当し、各スレッドが1回のAX+Y計算を実行する
- 729回の計算を実行するには729個のスレッドが必要であり、729回の並列計算が可能となる
- この時点ではプログラムはスレッド数とメモリリクエストの制約を受ける
GPUのキャッシュ機構
GPUの動作過程において、メモリ遅延、メモリ転送、メモリ帯域幅などのメモリ関連問題を可能な限り削減することが重要であり、その中でキャッシュメモリは特に重要である。NVIDIA Ampere A100のメモリ構造において、HBM Memoryのサイズは80Gであり、これがA100のビデオメモリサイズとなる。
レジスタファイルも一種のキャッシュメモリとして扱うことができる。レジスタはSM（Streaming Multiprocessor：ストリーミングマルチプロセッサ）実行ユニットに近接配置されており、実行ユニットからデータを迅速に取得できると同時に、L1キャッシュからデータを読み取ることも容易である。さらに、L2キャッシュはHBM Memoryにより近接配置されており、GPUが大量のデータをキャッシュに直接搬运することを可能にしている。これらの目標を同時に実現するため、GPUはマルチレベルキャッシュ设计了多级缓存。
80Gのビデオメモリは高い帯域幅を持つメモリであり、L2キャッシュサイズは40MB、すべてのSMが同一个L2缓存を共有する。L1キャッシュサイズは192kBであり、各SMが独立したキャッシュを持つ。同様に、各SMは独立したレジスタを持ち、各レジスタサイズは256KBである。SMの総数が108個であるため、レジスタの総サイズは27MB、L1キャッシュの総サイズは20MBとなる。
GPUとCPUのメモリ帯域幅と遅延を比較すると、GPUでは主メモリ（HBM Memory）をメモリ帯域幅（B/W：bandwidth）の基本単位とした場合、L2キャッシュの帯域幅は主メモリの3倍、L1キャッシュの帯域幅は主メモリの13倍となる。実際の計算時には、キャッシュされたデータを可能な限り早く使用し终了させてから次のデータを読み取る，这时会遇到延迟（Lentency）の問題。L1キャッシュの遅延を基本単位とした場合、L2キャッシュの遅延はL1の5倍、HBMの遅延はL1の15倍となる。したがって、GPUには独立したビデオメモリが必要である。
假设使用CPU将DRAM（Dynamic Random Access Memory）中的数据传输到GPU中进行计算，较高的延迟（25倍）会导致数据传输速度远低于计算速度，因此需要GPU拥有自己的高带宽内存HBM（High Bandwidth Memory）。GPU和CPU之间的通信和数据传输主要通过PCIe进行。
DRAM（Dynamic Random Access Memory）
DRAMはコンピュータメモリの基本类型であり、中央处理器（CPU）が快速アクセスするために、计算机程序とデータを一時的に存储するために使用される。静态随机存取存储器（SRAM）と比较すると、より高い存储密度と较低のコストを持つが、速度は遅い。这是计算机系统中最常用的内存类型之一，用于存储操作系统、应用程序和用户数据等内容。
DRAMの各存储单元由电容と晶体管组成，电容负责存储数据位（0または1），晶体管用于读取和刷新数据。由于电容会逐渐失去电荷，因此需要定期刷新（称为刷新操作）以保持数据的正确性，这也是称为"动态"的原因。DRAMは数据和程序の一時存储を提供し、较 빠른访问速度と相对的に低いコストを提供する。
| 存储类型 | 结构 | 工作原理 | 性能 | 应用 |
|---|---|---|---|---|
| DRAM（Dynamic Random Access Memory） | 一种基本的内存技术，通常以单层平面的方式组织，存储芯片分布在一个平面上 | 当读取数据时，电荷被传递到输出线路，然后被刷新。当写入数据时，电荷被存储在电容中。由于电容会逐渐失去电荷，因此需要周期性刷新来保持数据 | 具有较高的密度和相对较低的成本，但带宽和延迟相对较高 | 常用于个人电脑、笔记本电脑和普通服务器等一般计算设备中 |
| GDDR（Graphics Double Data Rate） | 专门为图形处理器设计的内存技术，具有较高的带宽和性能 | 在数据传输速度和带宽方面优于传统的 DRAM，适用于图形渲染和视频处理等需要大量数据传输的应用 | GDDR 与标准 DDR SDRAM 类似，但在设计上进行了优化以提供更高的数据传输速度。它采用双倍数据速率传输，即在每个时钟周期传输两次数据，提高了数据传输效率 | 主要用于高性能图形处理器（GPU）和游戏主机等需要高带宽内存的设备中 |
| HBM（High Bandwidth Memory） | 使用堆叠设计，将多个 DRAM 存储芯片堆叠在一起，形成三维结构 | 堆叠设计允许更短的数据传输路径和更高的带宽，同时减少了功耗和延迟。每个存储芯片通过硅间连接（Through Silicon Via，TSV）与其他存储芯片通信，实现高效的数据传输 | 具有非常高的带宽和较低的延迟，适用于高性能计算和 AI 等需要大量数据传输的领域 | 主要用于高端图形处理器（GPU）、高性能计算系统和服务器等需要高带宽内存的设备中 |
不同存储位置の帯域幅と計算強度を比較すると、HBMの計算強度を100とした場合、L2キャッシュの計算強度は39であり、每个数据只需要执行39个操作。L1缓存更少，计算强度只需要8个操作，这时候对于硬件来说非常容易实现。这就是为什么L1缓存、L2缓存和寄存器对GPU来说如此重要。可以把数据放在L1缓存里面然后对数据进行8个操作，使得计算达到饱和的状态，使GPU里面SM的算力利用率更高。
但是PCIeの带宽很低，整体的延迟很高，这将导致整体的算力强度很高，算力利用率很低。
| Data Location | Bandwidth(GB/sec) | Compute Intensity | Latency(ns) | Threads Required |
|---|---|---|---|---|
| L1 Cache | 19,400 | 8 | 27 | 32,738 |
| L2 Cache | 4,000 | 39 | 150 | 37,500 |
| HBM | 1,555 | 100 | 404 | 39,264 |
| NVLink | 300 | 520 | 700 | 13,125 |
| PCIe | 25 | 6,240 | 1,470 | 2,297 |
在带宽增加的同时线程的数量或者线程的请求数也需要相对应的增加，这个时候才能够处理并行的操作，每个线程执行一个对应的数据才能够把算力利用率提升上去，只有线程数足够多才能够让整个系统的内存处于忙碌的状态，让计算也处于忙碌的状态，因此看到GPU里面的线程数非常多。
GPUのスレッド原理
GPUの全体アーキテクチャと単一のSM（Streaming Multiprocessor）のアーキテクチャについて説明する。SMは基本的な演算単位としてみなすことができ、GPUは1つのクロックサイクルで複数のWarpを実行できる。1つのSMには64個のWarpが含まれており、4つのWarpは单独で并发的に実行可能である。GPU的设计者主要是通过增加线程和增加Warp来解决或掩盖延迟问题，而不是去减少延迟时间。
为了处理更多的计算任务，GPU SMs选择线程超配，每个SM共有2048个线程，整个A100有20多万个线程可供程序使用。在实际场景中程序无法使用所有线程，因此有些线程处于计算过程中，有些线程负责搬运数据，还有一些线程在同步等待下一次计算。很多时候会看到GPU的算力利用率不是非常高，但完全不觉得它慢是因为线程是超配的，远远超出大部分应用程序的使用范围，线程可以在不同的Warp上进行调度。
| Pre SM | A100 | |
|---|---|---|
| Total Threads | 2,048 | 221,184 |
| Total Warps | 64 | 6,912 |
| Active Warps | 4 | 432 |
| Waiting Warps | 60 | 6,480 |
| Active Threads | 128 | 13,824 |
| Waiting Threads | 1,920 | 207,360 |
本稿を通じて、GPUの基本的な動作原理から並列処理のアーキテクチャ、キャッシュ機構、そしてスレッド管理に至るまで、GPUの包括的な理解を深めることができた。GPUは深層学習や科学的シミュレーションなど、大規模な計算を要する分野で不可欠なハードウェアとなっており、その設計思想を理解することは、システム最適化やアルゴリズムの実装において重要な基盤となる。