GPUがメモリを読み取るときに何が起こるか

以前の投稿では、ベクトル加算カーネル — c[i] = a[i] + b[i]、floatごとに1つのスレッド — をnvccからワープまで追跡しました。カーネルがどのように起動されたかについては多くの詳細を説明しましたが、多くのことを省略しました。 今回は、省略した部分を補い、重要なSASS命令（グローバルロード）がハードウェアを通過する経路を追跡します。この場合、私のデスクの下にあるRTX 4090です。 私たちは、少なくとも原則として、パフォーマンス上の理由からこのようなリバースエンジニアリングを行っています（優れた理由については、Citadelマイクロベンチマーク論文の「なぜこれらの詳細が重要なのか」を参照してください）。よりプロダクションに関連性の高いGPUに適用された同じ作業については、このスペースにご注目ください。 この経路の詳細の多くはNVIDIAによって文書化されていません。少なくとも、私たちが望むレベルでは。そのため、ハードウェア自体でタイミング実験を実行して決定します。 調査しているCUDAカーネルは、その関数本体に2行あります。 __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を調べると、それらの行を駆動する命令が表示されます。 /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i] これらは、ベクトルbの要素をグローバルメモリからレジスタにロードするために使用されます。そこで、それらはaの要素に加算されてカーネルを実行します。 1つのLDG.Eは32レーンで4バイトを要求します。それを処理するには、4つの32バイトセクター、1つのキャッシュライン、1つのアドレス変換、クロスバーの通過、36個のL2スライスの一つ、そしてどこでもミスした場合、DRAMチップでのアクティベートと4回のカラムリードが必要です。 この命令がハードウェアを通過し、戻ってくるまでの旅を追っていきます。 状況設定：私たちのワープは、他の11個の常駐ワープと共に、SMの4つのサブパーティションの1つに存在します。各サイクル、サブパーティションのスケジューラは、利用可能なワープの1つを選択し、その次の命令を一度に32レーン全体に発行します。 私たちのワープは2回勝ちます。1回はIMAD.WIDEのため、そして数サイクル後（アドレスは現在R4とR5にあります）、LDGのためです。 私たちの物語はLDGから始まります。 ワープコーレッサー L115 nsTLBクロスバー L2127 nsコントローラー DRAM255 nsSMダイボード ワープからL1キャッシュへ 命令から始めましょう。 LDG.E R4, [R4.64] は、レジスタR4とR5に格納された64ビットアドレスから32ビットをグローバルロードします（R5は.64アノテーションのため表示されます。レジスタは32ビットサイズです）。結果をレジスタR4に格納します。 データをロードするには、まずそのアドレスをレジスタから取得する必要があります。 レジスタファイルの一行は、一度に32レーンすべてのR4を保持します（読み取りはまずオペランドコレクターでステージングされます。ステージングは、ソースがレジスタファイルの同じバンクを共有する命令のためにあります。バンクはサイクルごとに1回の読み取りをサービスします。2つのバンクがあり、レジスタ番号の最下位ビットで選択されるため、隣接するペアは常に両方をまたぎます）。もう一行はR5を保持します。 ワープは両方のエントリを読み取り、32個の異なる64ビットアドレスとして256バイトを読み取ります。これはレーンごとに1つのアドレスです。 レジスタ取得のコスト アドレスの読み取りは最大1サイクルかかります。レジスタからアドレスを取得する共有メモリロードは、発行から最初の使用まで24サイクルかかります。アドレスが即値である同じロードは23サイクルかかります（LDGは即値を取ることができません）。 すべてのアドレスが解決されると、命令はロード/ストアユニット（LSU）に発行されます。 LSUは命令とそのオペランドアドレスを受け取り、いくつかのアドレス演算（必要であれば）を実行します（このユニットは即値オフセットを追加できます。[R4.64]には追加するオフセットがありません）。そして、スコープロード（LDGはグローバルウィンドウを直接名前付けします）を行い、オペコード（バイナリで「これらのアドレスをロード」）、アクティブレーンの32ビットマスク、計算されたアドレス、そして結果を格納するレジスタ番号を送信します。 次の宛先はコーレッサーです。 各レーンの各LDG.E命令は4バイトを要求しますが、次の宛先であるL1キャッシュは32バイトセクター単位でアドレス指定されます。 コーレッサーの仕事は、4バイトのリクエストを処理するために必要な最小限のL1セクター数を把握することです。 コーレッサーは、ワープが要求した128バイトのために、4つの連続したセクターリクエストを発行すべきだと判断します。 1. L1キャッシュへの進入 ワープコーレッサー L115 nsTLBクロスバー L2127 nsコントローラー DRAM255 nsSMダイボード 4つの連続した32バイトセクターのリクエストは、L1キャッシュに送信されます。 L1キャッシュの組織単位はさらに細かく、128バイトラインです。私たちの4つの連続したセクターは、単一ラインの4つの部分を表すため、そのキャッシュラインへのリクエストがL1に対して行われます。 まず、そのラインがすでにキャッシュにあるかどうかを確認する必要があります。 キャッシュは、セットと呼ばれるグループに分割されています（技術的には、4090のL1キャッシュは4ウェイセットアソシアティブです。キャッシュは、完全にアソシアティブ（任意のキャッシュラインをキャッシュ内のどこにでも格納できる）と「ダイレクトマップ」（各キャッシュラインは1つの場所にしか格納できない）の連続体上にあります）。ラインのアドレスは、それが属するセットを決定します。 このカードのセットは4つのスロットを保持し、それぞれがラインを識別するタグを運びます。ルックアップは、要求するラインのタグと比較して、4つすべてを比較します。 使用されるアドレスは、プログラムで使用される仮想アドレスです（おそらく、L1にヒットするための変換コストを支払う必要がないようにするためです）。 ラインが配置されるセットは、ハッシュスキームによってラインの仮想アドレスから生成されます（これは単なるビットのスライスではなく、複雑なパリティスキームです。これにより、2のべき乗のストライドアクセス（行列の列、テンソルなどを考える）が同じセットにヒットし続けてチャーンするのを防ぎます）。 4つのタグのいずれかが一致し、要求するセクターがそのスロットにある場合、データが読み出され、ロードが完了します。 私たちのデータは初めてロードされるため、リクエストはミスし、メモリシステムをさらに下位に降りる必要があります。 L1ヒットのコスト L1ヒットは約15.4 ns — 40サイクルで返されます。この数値は、1つのスレッドがL1に常駐するラインのランダムな順列を介して依存チェーンを追いかけるレイテンシチェイスから得られます。 L2を探す：変換 ワープコーレッサー L115 nsTLBクロスバー L2127 nsコントローラー DRAM255 nsSMダイボード 仮想メモリは、プログラムが名前を付けるアドレスと、ハードウェアがデータを格納するアドレスの間に、1レベルの間接参照を置きます。 プログラムは独自の連続した空間を取得し、ハードウェアはその空間を物理ページに好きなように配置します。 変換は、それらの間のマップです。 L1は仮想アドレス指定されていたため、変換を気にする必要はありませんでした。ここからは、ハードウェアの言語を話す必要があります — L1ミスはSMを出る前に変換される必要があります。 物理アドレスと仮想アドレスの実際のマッピングは、ドライバでの割り当て時に確立されます。bが割り当てられたとき、ドライバはそれのために物理（2MiB）ページを選択し、割り当てを記録するページテーブルをVRAMに書き込みました。 変換ユニットは、仮想アドレスを受け取り、それらのテーブルに従って物理アドレスを返します。 SMは、最も最近の16個の変換を、ワープ間で共有されるTLBに保持します。 最初のロードはこのTLBでミスします。 変換のコスト プローブのいずれにおいても、TLBへのヒットコストは見られません。 ミスは、約4.4 ns — 11サイクルかかります。 同じリフィルコストは、このチップがマッピングできるすべてのページ、および任意のSMから0.1 ns以内で保持されるため、変換キャッシュの次のレベルは普遍的で非常に安価です。 変換が実行されると、128バイトラインごとに1つのリクエストが出力されます。これには、ラインの物理アドレスと、そこから必要なセクターのマスクが含まれます。私たちのリクエストは、すべての4つのセクターがマークされた単一のリクエストです。 リクエストはSMから出て、クロスバーを介してL2キャッシュに進みます。 L2で迷子に ワープコーレッサー L115 nsTLBクロスバー L2127 nsコントローラー DRAM255 nsSMダイボード リクエストはクロスバーを介して、36個の2 MiB L2スライスの1つに進みます。p