2026/08/22 1:16
GPU がメモリを読み込む際に起こること
RSS: https://news.ycombinator.com/rss
要約▶
Japanese Translation:
記事は、NVIDIA RTX 4090 におけるグローバルメモリアクセスの詳細な分析を提供し、ベクトル加算カーネルに対して顕著なパフォーマンスボトルネックを明らかにする。
LDG.E本文
RTX 4090 の CUDA コアにおけるメモリ経路追跡と SASS アナリシス：vadd
カーネル事例研究
vaddはじめに
前回の投稿では、ベクトル和演算
c[i] = a[i] + b[i]- 対象デバイス: RTX 4090
- 目的: パフォーマンス向上のためのリバースエンジニアリング（原理的には正当）
- 調査手法: CUDA コードとコンパイル後の SASS コードに基づき、実際のハードウェア上で計測実験を行う
注: NVIDIA は公開詳細が限定されています。そのため、実機での計測実験によって経路の真の姿を解明します。
調査対象：vadd
カーネル
vadd以下の 2 行で構成される CUDA コア関数が調査対象です。
__global__ void vadd(const float* a, const float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; }
コンパイルされた SASS コードでは、以下の命令が見られます。
: 2 つの入力アドレスを計算（IMAD.WIDE
など）。&b[i]
: グローバルメモリーからデータをレジスタに読み込む（LDG.E
のデータロード）。b[i]
本記事では、ベクトル
abLDG.E全体の経路概要
1 回の
LDG.E- 処理単位: 4 セクター（各セクター 32 バイト）、すなわち1 つのキャッシュライン。
- 必要なハードウェア資源:
- L1 キャッシュライン
- アドレス翻訳（TLB）
- クロスバースロープ（Crossbar）
- L2 スライス（36 個のうちいずれか）
- DRAM チップ（アクティベーションと 4 つのコラム読み込み）
1. レジスタからアドレスを取得
命令の詳細
LDG.E R4, [R4.64]R4R5
の注釈があるため、実際の読み出しには R5 が使用されます（R4 はアドレス保持用）。.64
アドレス読み出しのコスト
レジスタファイルの一行には 32 チェーン分まとめて R4 が格納されています。読み込みは「オペランドコレクター」で待機（staged）されます。
- バンク共有対策: 同じバンクを共有する場合は、1 サイクルあたり 1 回しか処理できないため、スタージングが必要です。
- 利用されるバンク: R4 と R5 は異なるバンクに保持され、両方使用して合計**256 バイト（32 チェーン×8 バイト/チェーン）**をアドレスとして取得します。
コスト比較
- レジスタからアドレス読み出し: 最大 1 サイクル追加
- 共有メモリアクセス発行 → 初回使用まで: 約 24 サイクル（immediate 定数の場合は 23 サイクル）
LSU（ロード/ストアユニット）への渡り
アドレスが解決されると、命令は LSU に渡されます。LSU は以下の情報を伝達します。
- オペランドのアドレス
- opcode（読み込み指示）
- アクティブなチェーンのマスク（32 ビット）
- 計算されたアドレス
- 結果が入るレジスタ番号
LSU は immediate オフセットの加算や、スコープ読み込み（
LDG2. コアレスカーへの進出と L1 キャッシュへの到達
コアレスカー（Coalescer）の役割
各チェーンの
LDG.E- コアレスカーの仕事: ワープが要求した 4 バイト（128 バイト相当）を満たすために必要な最小限の L1 セクター数を特定します。
- 結果: 1 つのワープからの要求に対し、4 つの連続したセクターの要求が発生し、1 つのキャッシュラインとして処理されます。
L1 キャッシュへの問い合わせ
L1 はセット関連型（4 ウェイ）です。
- 検索: リクエストされたタグが、対応する 4 つのスロットの中で一致するか確認します。
- ヒット時: データを返し、完了。
- ミス時（初回読み込みの場合）: キャッシュラインが存在しないため、より下位のメモリステージへ降ります。
アドレス空間: L1 は仮想アドレスを使用しています。物理アドレスへの翻訳はここでは行われません。
L1 ヒット時のコスト
ランダムな順列での測定結果です（依存チェーンを追いかけた一スレッドからの数値）。
- ヒットレイテンシー: 約 15.4 ns（40 サイクル）
3. TLB 翻訳による物理アドレス化
L1 ミスが発生すると、仮想アドレスから物理アドレスへの変換が必要です。
翻訳ユニットと TLB
- TLB (Translation Lookaside Buffer): SM 単位で共有される最近の 16 エントリを保持します。
- 動作: ドライバーが割り当て時に確立したページテーブルに従って、仮想アドレスを物理アドレスに変換します。
TLB コスト
- TLB ヒット: 測定で確認されたすべてのプローブで見られるコストではありません（0.1 ns 以内）。
- TLB ミス（リフィル）: 約 4.4 ns（11 サイクル） かかります。
翻訳後、物理アドレス付きの単一の要求（4 セクターすべてを含む）が生成され、SM を離れてクロスバースロープを通ります。
4. L2 キャッシュへの探索とミス
L2 アーキテクチャ
- 構成: 36 つの 2 MiB スライス（L1 と同じ構造：1024 セット、16 ウェイセット関連）。
- アクセス経路: クロスバースロープから、物理アドレスに基づいてスライスに到達します。
- 動作: L1 と同様、ラインがキャッシュに含まれていればヒットしますが、初回読み込みはミスとなります。
L2 ミスの処理
L2 ミスが発生すると、要求はメモリーコントローラを経由して GDDR6X DRAM チップへ送られます。
L2 ヒット時のコスト: 約 127 ns（330 サイクル） (注：これはヒットした場合ですが、本記事のフローではミスが前提です)
5. DRAM での読み出しと帰路
メモリーコントローラとチップ
要求はメモリーコントローラに引き継がれ、以下の動作を行います。
- アクティベーション: 行を開く（高コスト）。
- コラムリード: 4 つのコラムからデータを読み出す（低コスト）。
DRAM の構造と信号処理
- セル構造: トランジスタの後ろにコンデンサがあります。電荷は漏れるため、定期的なリフレッシュが必要です。
- 読み出しプロセス:
- アクティベーションコマンドでワードラインを駆動し、コンデンサの電荷をセンスアンプへ流す。
- センスアンプで電荷を増幅してフルレイルビットに変換。
- データピン（PAM4 シンボル）でシリアル化して出力（16 ピン×8 シンボル）。
データの帰還
- DRAM からデシリアル化した PAM4 バーストは、クロックと共にコントローラへ戻ります。
- データは L2 スライスのラインに書き込まれ、クロスバースロープを通って元の SM へ返されます。
- 最後に L1 のスロットを埋め、結果がレジスタ
に書き込まれます。R4
依存関係: レジスタへの書き込みにより「依存関係バリアー」がクリアされ、ワープは再び実行可能になります。
総計のレイテンシ
今回の往復経路（L1 ミス → TLB ミス → L2 ミス → DRAM）のコストは以下の通りです。
- 合計時間: 約 255 ns
- サイクル数: 約 660 サイクル
この間、ワープはバリアーにより停止しています（チップの他の部分や SM は動作中）。単一の読み込みのレイテンシはノイズに埋もれます。
アペンドックス：プローブと反証実験手法
計測には以下の 2 つの儀器を使用しました。
- レイテンシーチェース: ラインを指し示すポインターサイクルを 20,000 回跳躍させて実行し、平均ナノ秒を測定。
- L1: 15.4 ns
- L2: 127.4 ns
- DRAM: 255.4 ns
- ハードウェアカウンター (ncu): グラディエントとして固定オーバーヘッドを相殺。
L2 スライスへの到達関数（推論）
L2 は物理的にインデックス化されているため、アドレス変換には複雑なパリティ関数が使用されています（4090 の 36 スライスをマッピング）。
def parity(x): return bin(x).count("1") & 1 def _state(a, N): wide = (N == 48) # L40S: 4 slices/controller b35 = (1 << 35) if wide else 0 # controller selection P1c = parity(a & 0x76A990400) P1 = parity(a & (0x76A990400 ^ b35)) P2 = parity(a & 0x2CCF7B000) A = ((a >> 15) + 2*parity(a & 0x3C9041000) + parity(a & (0x2882B0800 ^ b35)) + 2) % 3 # slice selection within controller g = ((a + (1 << 16)) >> 17) % 9 q0 = parity(a & 0x8000) q1 = parity(a & 0x5985E0500) q2 = parity(a & (0x2354E4400 ^ b35)) q3 = parity(a & 0x3C9041000) carry = 1 if q0 + q1 + q2 >= 2 else 0 start = (5 + 7*q0 + 5*q1 + 2*q2 + q3 - carry) % 9 o = (g - SHIFT[A] - start) % 9 Lf = 2 if (q0 ^ q1 ^ q2) == 0 else 1 return P1c, P1, P2, A, q2, o // 3, (1 if (o % 3) >= Lf else 0) def slice_of(a, N=36): # ... (controller calculation logic omitted for brevity) ... return controller * 3 + B
検証結果
予測したスライスからのライン抽出率による検証（Mload/s）：
- 1 スライス: 1.00×
- 2 スライス: 2.00×
- 4 スライス: 4.00×
- 36 スライス (全スライス): 68,085 (予測値の約 35 倍)
DRAM リフレッシュの影響
DRAM セルは電荷を漏らすため、定期的なリフレッシュが発生します。
- 影響: アクセスの約 2% が通常のレイテンシーより最大約 210 ns 高いハード ceilings に分散。
- 特徴: フォールトレスター（Stall）として固定長のシグネチャとして観測可能。
補足情報：特定の技術的注釈
以下の点について重要です。
- 仮想/物理アドレス: L1 はインデックスとタグの両方で仮想アドレスを使用しています。L2 は物理アドレスを使用します（翻訳は L2 前で発生）。
- セットインデックス: 8 ビットは固定サブセットの排他論理和（XOR）で構成されます。
- スロット制限: L2 は一部のセクターしか持っていないラインを含めることが可能です。
- ページテーブル:
ツールでホストから GPU のページテーブルを辿り、ディレクトリエントリーの volatile ビットを確認できます。nvdebug - TLB 構造: 全関連型（Associativity）、16 エントリ、SM 単位で共有（ワープ間共有）、LRU で置換。
- 行サイズ測定: 連続した読み込みの実験により、行サイズが 1 KiB (32 コラム) であることを確認しています。
- アクティベーションコスト: 新しい行を開く場合、約 15 倍のコストがかかります（既存行なら約 3.4 ns の追加）。