Tiny GPU: Verilogで実装した最小限のGPU
(github.com/adam-maj)- tiny-gpuは、GPUがハードウェアレベルでどのように動作するかを基礎から学ぶためのVerilogベースの最小GPU実装であり、グラフィックス専用ハードウェアよりもGPGPUやMLアクセラレータの共通原理に重点を置いている
- 実装は、15個未満の文書化されたVerilogファイル、アーキテクチャとISAドキュメント、行列加算・乗算カーネル、カーネルのシミュレーションと実行トレース対応で構成される
- GPUは一度に単一カーネルを実行し、プログラムメモリとデータメモリをロードして
thread_countを設定した後、start信号を立ててカーネルを開始する - 単純化のため、各コアは一度に1つのblockを処理し、各スレッドはALU・LSU・PC・レジスタファイルを持つが、すべてのスレッドが各命令の後に同じPCへ収束すると仮定している
- 現代GPUの多層キャッシュ、共有メモリ、メモリコアレッシング、パイプライニング、ワープスケジューリング、分岐ダイバージェンス、バリアといった機能の大半は除外し、学習向けの構造を優先している
tiny-gpuが解決しようとする問題
- CPUにはアーキテクチャから制御信号まで学べる資料が多い一方で、現代GPUの低レベルな技術的詳細は競争の激しい市場のため大半が独占的なままになっている
- GPUプログラミングに関する資料は多いが、GPUがハードウェアレベルでどのように動作するかを学べる資料はほとんどない
- オープンソースGPU実装であるMiaowとVeriGPUは、機能の完成度と動作を目標としているため構造が複雑である
- tiny-gpuは製品レベルのグラフィックスカードにある多くの複雑さを取り除き、現代のハードウェアアクセラレータに共通する中核要素へ集中している
- GPUアーキテクチャの重要な構成要素
- SIMDプログラミングモデルがハードウェアで実装される方式
- 限られたメモリ帯域幅をGPUが扱う方式
全体アーキテクチャ
- tiny-gpuは一度に1つのカーネルだけを実行するよう設計されている
- カーネル実行手順は次のとおり
- グローバルプログラムメモリにカーネルコードをロード
- データメモリに必要なデータをロード
- デバイス制御レジスタに実行するスレッド数を指定
start信号をhighに設定してカーネルを実行
- GPUは次のユニットで構成される
- デバイス制御レジスタ
- ディスパッチャ
- 可変個数のコンピュートコア
- データメモリとプログラムメモリ用のメモリコントローラ
- キャッシュ
カーネル実行とスレッド割り当て
- デバイス制御レジスタはカーネル実行メタデータを保存する役割を持ち、tiny-gpuでは実行する総スレッド数である
thread_countだけを保存する - ディスパッチャはカーネル開始時にスレッドを複数のコンピュートコアへ分配する
- 並列実行可能なスレッドの束をblockとして構成する
- 利用可能なコアへblockを送って処理する
- すべてのblockの処理が終わるとカーネル実行完了を通知する
- 単純化されたコアは一度に1つのblockを処理する
- 各スレッドは専用のALU、LSU、PC、レジスタファイルを持つ
- これらの資源上でスレッド命令実行を管理することが、GPUの難しい課題の1つである
メモリ構造とコントローラ
- GPUは外部グローバルメモリとインターフェースするよう作られており、単純化のためデータメモリとプログラムメモリを分離している
- データメモリ仕様
- 8ビットアドレッシング
- 合計256行
- 8ビットデータ
- 各行は256未満の値を保存
- プログラムメモリ仕様
- 8ビットアドレッシング
- 合計256行
- 16ビットデータ
- ISAに従って各命令は16ビット
- メモリコントローラはコアから来るメモリ要求を追跡し、実際の外部メモリ帯域幅に合わせて要求を制限し、応答を正しい資源へ渡す
- 各メモリコントローラはグローバルメモリ帯域幅に応じた固定数のチャネルを持つ
- キャッシュは作業中の機能であり、外部メモリから取得したデータをデバイス内SRAMに保存することで、その後の要求でより高速に取得でき、新しいデータにメモリ帯域幅を使えるようにする
コア内部構成
- 各コアにはスレッド実行を管理する単一のスケジューラがある
- tiny-gpuのスケジューラは1つのblockの命令を最後まで実行してから新しいblockを取得し、すべてのスレッド命令を同期された順序で実行する
- より高度なスケジューラでは、パイプライニングやワープスケジューリングによって資源利用率を高められる
- スケジューラの主な制約は、グローバルメモリからデータをロード・保存するときに発生するレイテンシである
- ほとんどの命令は同期的に実行できる
LDRやSTRのようなload-store操作は非同期であり、長い待ち時間を中心に命令実行を構成する必要がある
- Fetcherは現在のプログラムカウンタの命令をプログラムメモリから非同期に取得する
- Decoderは取得した命令をスレッド実行用の制御信号へデコードする
- 各スレッドのレジスタファイルは計算中のデータを保持し、SIMDパターンを可能にする
- 読み取り専用レジスタには
%blockIdx、%blockDim、%threadIdxが入っている - カーネルはローカルスレッドIDに応じて異なるデータで実行できる
- 読み取り専用レジスタには
- 各スレッドのALUは
ADD、SUB、MUL、DIV算術命令を処理する CMPは2つのレジスタ差分の結果が負・0・正のどれかを出力し、その結果をPCユニットのNZPレジスタに保存する- 各スレッドのLSUはグローバルデータメモリにアクセスし、
LDRとSTR、および非同期メモリ待ち時間を処理する - 各スレッドのPCは次に実行する命令を決定する
- デフォルトでは命令ごとに1ずつ増加する
BRnzpは直前のCMPが設定したNZPレジスタ条件が一致すると、特定のプログラムメモリ行へ分岐する- ループと条件文はこの方法で実装される
- tiny-gpuは単純化のため、すべてのスレッドが各命令の後に同じPCへ収束すると仮定している
- 実際のGPUでは個別スレッドが異なるPCへ分岐でき、このとき一緒に処理されていたスレッドグループが複数の実行フローへ分かれる分岐ダイバージェンスが発生する
ISA
- tiny-gpuは行列加算や行列乗算のような概念実証用の単純なカーネルを実行するため、11命令のISAを実装している
- 対応命令
BRnzp:NZP条件が一致したら別のプログラムメモリ行へジャンプCMP: 2つのレジスタ値を比較し、結果をNZPレジスタへ保存ADD,SUB,MUL,DIV: テンソル演算向けの基本算術演算LDR: グローバルメモリからデータをロードSTR: グローバルメモリへデータを保存CONST: 定数値をレジスタへロードRET: 現在のスレッド実行の終了を通知
- 各レジスタは4ビットで指定され、合計16個のレジスタを持つ
R0からR12までの13個は読み書き可能な汎用レジスタ- 最後の3個はSIMDに必要な
%blockIdx、%blockDim、%threadIdxを提供する読み取り専用特殊レジスタ
実行フロー
- 各コアは命令実行時に次の段階の制御フローに従う
FETCH: 現在のPCの次の命令を取得DECODE: 命令を制御信号へデコードREQUEST:LDRまたはSTRが必要な場合、グローバルメモリへデータを要求WAIT: 必要に応じてグローバルメモリ応答を待機EXECUTE: データに対する計算を実行UPDATE: レジスタファイルとNZPレジスタを更新
- この制御フローは単純さと理解しやすさのために構成されている
- 実際の実装では、一部段階を圧縮して処理時間を最適化したり、パイプライニングでコア資源上の複数命令実行を調整したりできる
- 各スレッドは専用レジスタファイル内のデータを対象に同じ実行経路に沿って計算する
- CPUのダイアグラムに似ているが、
%blockIdx、%blockDim、%threadIdxが読み取り専用レジスタにあることでSIMD機能を可能にしている点が異なる
例示カーネル
- ISAの概念実証として、行列加算と行列乗算カーネルが作成されている
- リポジトリのテストファイルでは、これらのカーネルをGPU上で完全にシミュレーションし、データメモリ状態と全実行トレースを生成できる
-
行列加算
matadd.asmは1 x 8行列2つを加算する- 8個の要素ごとの加算をそれぞれ別スレッドで実行する
%blockIdx、%blockDim、%threadIdxレジスタを使ってSIMDプログラミングを示すLDRとSTR命令を使い、非同期メモリ管理も含む
-
行列乗算
matmul.asmは2 x 2行列2つを乗算する- 対応する行と列の内積を要素ごとに計算する
CMPとBRnzpを使ってスレッド内部の分岐を示す- すべての分岐が再び収束するため、現在のtiny-gpu実装で動作する
シミュレーション
- カーネルシミュレーションを実行するにはiverilogとcocotbが必要
- 準備手順
brew install icarus-verilogとpip3 install cocotbでVerilogコンパイラとcocotbをインストール- sv2vの最新版をダウンロードして展開し、バイナリを
$PATHに追加 - リポジトリルートで
mkdir buildを実行
- カーネルシミュレーションは
make test_mataddとmake test_matmulで実行する - 実行結果は
test/logsのログファイルとして出力される- 初期データメモリ状態
- カーネルの全実行トレース
- 最終データメモリ状態
- 各ログファイルの冒頭には入力行列が表示され、末尾の最終データメモリには結果行列が表示される
- 実行トレースには各サイクルごとの全コア・全スレッドの実行状態が含まれる
- 現在の命令
- PC
- レジスタ値
- 状態情報
意図的に省かれている高度なGPU機能
- tiny-gpuは単純化のため、現代GPUの性能・機能向上要素の大半を除外している
-
多層キャッシュと共有メモリ
- 現代GPUはグローバルメモリアクセスを減らすために複数のキャッシュ階層を使う
- tiny-gpuは要求元資源とメモリコントローラの間に最近のデータを保存する単一のキャッシュ階層だけを実装している
- 多層キャッシュは頻繁に使うデータを使用箇所により近い場所へキャッシュしてロード時間を短縮する
- GPUは同じblock内のスレッドが共有結果をやり取りできるよう、共有メモリを使うこともある
-
メモリコアレッシング
- 並列実行中の複数スレッドは、行列の隣接要素のように連続したアドレスへアクセスすることが多い
- メモリコアレッシングはキューに積まれたメモリ要求を分析し、隣接する要求を1つのトランザクションへ統合する
- アドレッシングに使う時間を減らし、要求をまとめて処理することが目的である
-
パイプライニング
- tiny-gpuのコアは、あるスレッド束の1命令の実行が終わってから次の命令を開始する
- 現代GPUは依存関係のある命令の順次実行を保証しつつ、複数の逐次命令実行をストリーミングする
- 非同期メモリ要求待ちのような状況でもコア資源が遊ばないようにし、資源利用率を高める
-
ワープスケジューリング
- blockを一緒に実行可能なスレッドの束であるwarpに分ける
- あるwarpが待機中のとき、別のwarpの命令を実行して単一コアで複数warpを同時に扱う
- パイプライニングと似ているが、異なるスレッドの命令を扱う
-
分岐ダイバージェンス
- tiny-gpuは単一バッチ内のすべてのスレッドが各命令後に同じPCを持つと仮定している
- 実際にはデータに応じて個別スレッドが別々の行へ分岐できる
- 異なるPCを持つスレッドは別個の実行フローに分かれ、再収束する時点も管理しなければならない
-
同期とバリア
- 現代GPUでは、同じblock内のスレッドグループが特定地点に全員到達するまで待つようバリアを設定できる
- スレッド同士で共有データを交換する必要がある場合、データ処理が完了したことを保証するのに有用である
次の作業
- 今後の改善項目は次のとおり
- 簡単な命令キャッシュの追加
- Tiny Tapeout 7でGPUを使えるようにするアダプタの構築
- 基本的な分岐ダイバージェンスの追加
- 基本的なメモリコアレッシングの追加
- 基本的なパイプライニングの追加
- サイクル時間改善のための制御フローとレジスタ使用の最適化
- グラフィックス機能を示すための基本的なグラフィックスカーネルの作成、または単純なグラフィックスハードウェアの追加
- リポジトリの改善を望むユーザーはPRで貢献できる
1件のコメント
Hacker Newsのコメント
例外的にIntelはGPU技術文書を数多く公開している: https://kiwitree.net/~lina/intel-gfx-docs/prm/
i810/815のマニュアルもオンラインで見つかるし、855/910/915/945が抜けている965以前の奇妙な空白期間を除けば、文書化はかなり継続的に行われていた
現在および過去の製品の命令セットアーキテクチャ文書まで含まれているが、関心のある愛好家向けの高レベルな説明というより、実装者向けの文書に近いように見える
すべての企業がこのやり方に従ってくれればと思う
ただ、これはSIMD補助プロセッサに近いと思う
GPUと呼ぶには、少なくとも何らかの形のディスプレイ出力は必要だと考えている
最近はNvidiaなどがサーバー専用のグラフィックアーキテクチャ派生品もGPUとして売っていて、用語がかなり緩くなっているのは分かるが、GPU設計においてグラフィックス部分はいまでも複雑さのかなりの部分を占めている
出力しないGPUにも依然として用途はある
私の職場には中級クラスのQuadroを積んだワークステーションが75台ほどあるが、カードにはmini-DisplayPortしかなく、会社がHDMIケーブルしか買ってくれないので、すべて内蔵グラフィックスにつながれている
それでもそれらのカードはソフトウェアを高速化し、グラフィックスを処理しており、単に画面出力をしていないだけだ
別の例もある: https://github.com/jbush001/NyuziProcessor
TSMCや他のファウンドリでこの種のプロセッサを経済的に生産するには、どの程度の数量が必要なのだろうか?
FPGAをやってみたいのだが、正直どこから始めればいいのか見当もつかず、分野全体がかなり威圧的に感じられる
最終目標はLLM向けのアクセラレータカードを作ることだが、完全に思いつきで決めた目標とはいえ、このプロジェクトと重なる部分は多そうで、おそらくより大きなモデルをロードするためのメモリオフロード部分くらいしか違わない気がする
FPGA入門は複数の下位技術に分解して考えるべきで、期待値も調整しないといけない
ソフトウェアエンジニアに対して、最初から原理からコンピュータ全体を作り、命令セットアーキテクチャを書き、機械語を理解し、それをアセンブリに変換し、Pythonコードでアプリケーションを作るためにプログラミング言語まで開発しろとは普通期待しない
上から始めてスタックを下っていくのが正しい
複雑さを抽象化し、あらかじめ用意されたIPでシステムを構築することに集中すれば、FPGA設計はかなり簡単になる
普通はMATLABのようなものを勧める。リファレンスデザインのあるDevKitでHDL Coderを使って初期アプリケーションを作れるからだ
そうでなければ、デジタルコンピューティングアーキテクチャ、Verilog、タイミング、トランシーバー/入出力、ピン配置、Quartus/Vivado、シミュレーション/検証、組み込みシステムなどを学ぶという膨大な負担が生じる
要するに、システムレベル設計から始めて、プラグアンドプレイのIPを持ってきてトップレベルで接続する方法を学び、そのモジュールをあらかじめ作られた基準設計に入れてみればよい
その後、徐々に層をはがしながら下の複雑さを明らかにしていけばよい
Harris&Harrisを読みながら演習を並行して進めることもできる
演習と本を終えたら、自分のプロジェクトを始める段階になる
ちなみにHackerMojoでは毎週の集まりもあり、ValleyにいなくてもZoomで参加できる
そこから始めればよい
ハードウェアは実際には簡単な部分でもあり、製造面では難しい部分でもある
alwaysブロックで non-blocking 代入 と blocking 代入 演算子を混在させて使う理由はあるのか？opencores という、複数のオープンソースHDLプロジェクトが集まっていたサイトがあった
最近、HPCレベルの大規模分散HDLシミュレータで良いものがあるのか気になる
RTLレベルのシミュレーションに現代のGPUを活用するのは妥当に思える
同じサイトではなく、似た別の場所なのか？
現代のCUDAコアのようなものでは、実際に命令として除算を持つのが普通なのか、それとも普通はソフトウェアでエミュレーションするのか？
実際のハードウェア除算回路は非常に面積を食うので、GPU ALUに入るとは予想していなかった
Verilogでは
DIV: begin alu_out_reg <= rs / rt; endと1行書くのはあまりにも簡単だが、その1行はシリコンをかなり消費する単にVerilogだけをシミュレーションしていると、その事実が見えないこともある
プロジェクトはシミュレーションで止まっていて、実際のハードウェアとして作るにははるかに多くの作業が必要だ
個人的には、こういうものは別の名前で呼ぶべきだと思う
両者の隔たりは縮まりつつあり、どちらも互いの機能を取り込みつつあるが、それでもなおかなりの違いがある
私は アムダールの法則 と関係があると思う [0]
その意味では、CPUはレイテンシ最適化プロセッサ、GPUはスループット最適化プロセッサと呼べる
さらに具体的には、[1] CPUを長く深いデータ依存性のプロセッサ、GPUを広く浅いデータ依存性のプロセッサと呼ぶこともできる
[0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
[1]: https://en.wikipedia.org/wiki/Data_dependency
テンソルは単なるn次元配列だ
その上にソフトウェアやファームウェアを載せて、GPUのように動作させることもできる
EDK2のソースからつなぎ合わせようとしたが、どこまでがQEMU特化なのか不明だ
tiny-gpuが、すべてのスレッドが各命令の後に同じプログラムカウンタへ「収束」すると仮定しているのは、あまりにも素朴な単純化だ実際のGPUでは個々のスレッドが異なるPCへ分岐でき、最初は一緒に処理されていたスレッド群が別々の実行へ分かれる 分岐発散 が生じる
シリコンでGPUを作る前に、まずGPUプログラミングをやってみたほうがよかっただろう
しかもSIMDと呼ぶにしても何かしっくりこない
この人は以前、他人の回路をつなぎ合わせてLEDを点滅させ、CPUを作ったと言っていたあの人だ
__syncthreads()を呼ぶのと同じようなものではないか？