GPUチップの設計方法
Chapter 5 Gpu Memory System Design

第5章: GPUメモリシステムの設計

グラフィックス処理ユニット(GPU)は、幅広い用途において高パフォーマンスと高エネルギー効率を達成できる、高度に並列化された、プログラマブルなアクセラレータに進化してきました。メモリシステムは、多数の並列スレッドにデータへの高速アクセスを提供しなければならないため、現代のGPUアーキテクチャにとって重要な構成要素です。本章では、GPUメモリシステムの設計における主要な要素について探ります。これには、GPUで使用されるDRAM技術、メモリコントローラとアービトレーション、共有メモリとキャッシュ、効率的なメモリ利用のためのテクニックなどが含まれます。

GPUのためのDRAM技術

ダイナミックランダムアクセスメモリ(DRAM)は、GPUを含む現代のコンピューティングシステムにおける主メモリの実装に使用される主要な技術です。DRAMは高密度と比較的低コストを提供しますが、キャッシュやレジスタファイルのようなオンチップメモリに比べて、アクセスレイテンシが高く、帯域幅が低いという特性があります。

GPUは一般的に、低レイテンシよりも高帯域幅を重視した専用のDRAM技術を採用しています。GPUで使用される一般的なDRAM技術には以下のようなものがあります:

  1. GDDR (Graphics Double Data Rate): GDDRは、グラフィックスカードやゲームコンソールに最適化された専用のDRAM技術です。標準のDDRDRAMよりも広いバスと高クロック速度を使用することで、より高い帯域幅を提供します。GDDR5とGDDR6が最新のバージョンで、それぞれ最大512 GB/sと768 GB/sの帯域幅を提供します。

  2. HBM (High Bandwidth Memory): HBMは、非常に高い帯域幅と低消費電力を提供する高性能の3D積層DRAM技術です。HBMは複数のDRAMダイを積層し、シリコン貫通接続(TSV)を使用して相互接続することで、従来のDRAMよりも格段に高いデータ転送レートを実現しています。HBM2は最大1 TB/sの帯域幅を提供できます。

図5.1は、従来のGDDRメモリと3D積層HBMの違いを示しています。

   GDDRメモリ
```HBM メモリ
  ____________                   ______________________
 |            |                 |  ___________________  |
 |   DRAM     |                 | |                   | |
 |   Chips    |                 | |      DRAM Dies    | |
 |            |                 | |___________________| |
 |            |                 |           .          |
 |            |                 |           .          |
 |            |                 |           .          |
 |____________|                 |  ___________________  |
      |                         | |                   | |
     PCB                        | |  Logic Die (GPU)  | |
                                | |___________________| |
                                |______________________|

図5.1: GDDR メモリと HBM メモリのアーキテクチャの比較

GPU の電力予算、フォームファクター、ターゲットアプリケーションなどの具体的な要件に応じて、DRAM テクノロジーが選択されます。ゲーミングや専門グラフィックス用の高性能 GPU では GDDR6 が、データセンターや HPC GPU では電力効率が重要な HBM2 がよく使われます。

## メモリコントローラとアービトレーション

メモリコントローラは、GPU とオフチップ DRAM の間のデータの流れを管理する責任を持っています。GPU コアからのメモリリクエストを処理し、DRAM コマンドをスケジューリングし、帯域幅の活用を最大化し、レイテンシを最小化するようなメモリアクセスパターンを最適化します。

GPU のメモリコントローラは通常、高い帯域幅と並列アクセスを提供するマルチチャネル設計を採用しています。各メモリチャネルは1つ以上の DRAM チップに接続され、それぞれコマンドバスとデータバスを持っています。メモリコントローラはメモリリクエストをavailable なチャネルに分散して、並列性を最大化し、チャネル競合を避けます。

図5.2は4チャネルのGPUメモリコントローラの簡略図を示しています。

GPU コア | | | | | メモリメモリ コントローラ | | | | Ch0 Ch1 Ch2 Ch3 | | | | DRAM DRAM DRAM DRAM

図5.2: 4チャンネルのGPUメモリコントローラ

メモリアービトレーションは、複数の未処理リクエストがある場合に、どのメモリリクエストを最初に処理するかを決める過程です。GPUは、メモリシステムのパフォーマンスと公平性を最適化するために、さまざまなアービトレーションポリシーを採用しています:

1. **先着順(FCFS)**: 最も単純なアービトレーションポリシーで、リクエストが到着した順番で処理されます。FCFSは公平ですが、リクエストの並び替えが行われないため、パフォーマンスが低下する可能性があります。

2. **ラウンドロビン(RR)**: リクエストが循環順に処理されるため、すべての要求者に同等の優先度が与えられます。RRは公平性を提供しますが、ロケーリティや緊急性を最適化できない可能性があります。

3. **優先度ベース**: さまざまな基準に基づいて、リクエストに優先度が割り当てられます。たとえば、リクエストのタイプ(読み取りや書き込みなど)、ソース(テクスチャやL2キャッシュなど)、リクエストの経過時間などです。高い優先度のリクエストが先に処理されます。

4. **デッドラインアウェア**: リクエストのデッドラインに基づいて、スケジューリングが行われます。これは、リアルタイムグラフィックスアプリケーションにとって特に重要です。

5. **ロケーリティアウェア**: メモリコントローラは、近接するメモリ領域へのリクエストをまとめて処理することで、ローバッファヒットを最大化し、DRAM プリチャージとアクティベーションのオーバーヘッドを最小化しようとします。

高度なGPUメモリコントローラは、しばしばこれらのアービトレーションポリシーの組み合わせを使用して、パフォーマンス、公平性、リアルタイム要件のベストバランスを実現しています。

## シェアードメモリとキャッシュ

GPUは、メインメモリに対するレイテンシとバンド幅の要求を低減するために、ソフトウェア管理とハードウェア管理の両方のキャッシュを含む階層的メモリシステムを採用しています。

### シェアードメモリ

シェアードメモリは、ソフトウェア管理の、スレッドブロック(NVIDIA)または作業グループ(AMD)内のスレッド間で共有される、オンチップのメモリ領域です。以下为Japanese翻訳版本:

kgroup (OpenCL)は、スレッドブロック内のデータ移動と再利用を明示的に管理できるようにするユーザー制御型キャッシュです。

共有メモリは通常、低待ち時間と高帯域幅アクセスを提供するために、高速な多ポート SRAMバンクを使って実装されています。各バンクは1サイクルに1つのメモリリクエストをサービスできるため、ハードウェアは同じバンクへの競合するアクセスを回避するために仲裁する必要があります。

図5.3は、GPUコアの共有メモリの構成を示しています。

スレッドブロック

| _________________ | | | スレッド 0 | | | || | | . | | . | | . | | _________________ | | | スレッド N-1 | | | || | |_______________| | | | | | 共有メモリ | | ____________ | | | バンク 0 | | | |____| | | | バンク 1 | | | || | | . | | . | | . | | | バンクM-1 | | | |__________| | ||

図5.3: GPUコアの共有メモリ構成

共有メモリの適切な使用は、DRAMへのアクセス回数を減らすことで、GPUカーネルのパフォーマンスを大幅に向上させることができます。ただし、効率的なデータ共有とバンクコンフリクトの回避のために、慎重なプログラミングが必要です。

### ハードウェア管理キャッシュ

ソフトウェア管理の共有メモリに加えて、GPUはデータローカリティを自動的に活用し、DRAM アクセスを削減するためにハードウェア管理のキャッシュも採用しています。GPUで一般的なハードウェア管理キャッシュのタイプは以下の通りです:

1. **L1データキャッシュ**: 最近アクセスされたグローバルメモリデータを格納する小型のコア専用キャッシュ。L1キャッシュはグローバルメモリアクセスの待ち時間を削減するために使用されます。

2. **テクスチャキャッシュ**: 読み取り専用テクスチャデータへのアクセスを最適化するための専用キャッシュ。こちらが日本語訳版のマークダウンファイルになります。コードの部分は翻訳せず、コメントの部分のみ翻訳しました。

 GPU テクスチャデータ  
テクスチャキャッシュは2D空間的局所性を最適化しており、ハードウェアアクセラレーションによるフィルタリングと補間操作をサポートしています。

3. **定数キャッシュ**: 頻繁にアクセスされる定数データを格納する小さな読み取り専用のキャッシュ。定数キャッシュはワープ内のすべてのスレッドにブロードキャストされるため、多くのスレッドで共有されるデータに効率的です。

4. **L2キャッシュ**: GPUコアと主メモリの間にある大きな共有キャッシュ。L2キャッシュはL1キャッシュから追い出されたデータを格納し、DRAMアクセスの数を減らすために使用されます。

図5.4は、ハードウェアで管理されるキャッシュを備えた典型的なGPUメモリ階層を示しています。

GPUコア 0 GPUコア 1 GPUコア N-1

| | | | | | | L1 Data | | L1 Data | | L1 Data | | キャッシュ | | キャッシュ | | キャッシュ | || || || | | | | | | | テクスチャ | | テクスチャ | | テクスチャ | | キャッシュ | | キャッシュ | | キャッシュ | || || || | | | | | | | 定数 | | 定数 | | 定数 | | キャッシュ | | キャッシュ | | キャッシュ | || || |______________| | | | |_________________|_________________| | | | | | L2 キャッシュ| |_____________| | | メインメモリ

図5.4: ハードウェアで管理されるキャッシュを持つGPUメモリ階層メモリ管理されたキャッシュは、データのローカリティを自動的に活用し、DRAM アクセスの回数を減らすことで、GPU アプリケーションのパフォーマンスを向上させるのに役立ちます。しかし、CUDA やOpenCLのような並列プログラミングモデルの文脈では、キャッシュの整合性と一貫性の課題も引き起こす可能性があります。

## 効率的なメモリ利用のためのテクニック

GPU メモリシステムの効率的な利用は、高パフォーマンスと省エネルギー効率を達成するために非常に重要です。GPU アプリケーションでメモリ使用を最適化するための主なテクニックには以下のようなものがあります:

1. **コアレスシング**: ワープ内のスレッドからのメモリアクセスを隣接するメモリ領域に配置することで、ハードウェアがそれらを単一の幅の広いメモリトランザクションにまとめることができます。コアレスシングは DRAM バンド幅の利用率を最大化し、メモリトランザクションの数を削減します。

2. **データレイアウトの最適化**: メモリ上でデータ構造を組織化し、空間的ローカリティを最大化し、キャッシュミスを最小限に抑えます。これには、同じタイプのデータ要素をまとめるストラクチャオブアレイ (SoA) レイアウトや、同じ構造に属するデータ要素を一緒に保持するアレイオブストラクチャ (AoS) レイアウトなどの手法が含まれます。

3. **キャッシングとプリフェッチング**: メモリアクセスパターンの時間的および空間的ローカリティを活用することで、ハードウェアマネージドのキャッシュを効果的に活用します。これには、データタイリングによってデータを小さな塊に分割してキャッシュに収まるようにしたり、ソフトウェアプリフェッチによって必要になる前にデータをキャッシュに読み込んだりするテクニックが含まれます。

4. **メモリアクセススケジューリング**: ロウバッファのヒットを最大化し、DRAM のプリチャージやアクティベーションのオーバーヘッドを最小限に抑えるためにメモリアクセスを並べ替えます。これは、メモリコントローラーのハードウェアメカニズムや、アクセスパターンの最適化やデータレイアウトの変換などのソフトウェアテクニックによって実現できます。

5. **圧縮**: データ圧縮テクニックを適用して、メモリとGPUコア間で転送されるデータのサイズを縮小します。これにより、メモリバンド幅の使用率が向上し、エネルギー効率も改善される可能性があります。ネットワークの帯域幅のボトルネックを緩和し、データ移動に関連するエネルギー消費を削減する。

6. **メモリ仮想化**: GPUアプリケーションの統一された連続アドレス空間を提供するために、バーチャルメモリの技術を活用する。これにより、メモリ管理が柔軟になり、需要ページングなどの機能が利用できるようになり、メモリフットプリントの削減とシステムの利用率の向上に役立ちます。

図5.5は、GPUメモリシステムのコンテキストでこれらの技術の一部を示しています。

GPU コア | | | | | コアレッシング| || | | | | | データレイアウト| | 最適化 | || | | | | | キャッシングと | | プリフェッチング| || | | | | | メモリ | | アクセス | | スケジューリング| || | | | | | 圧縮 | || | | | | | メモリ | |仮想化 | || | DRAM

図5.5: GPUメモリシステムでの効率的なメモリ利用のための手法

1. **コアレッシング**: ワープ内のスレッドからのメモリアクセスを隣接するメモリ位置に配置し、ハードウェアがそれらを単一の幅広いトランザクションに結合できるようにする。コアレッシングは、DRAMの帯域幅の利用を最大化し、メモリトランザクションの数を減らします。

   例:
   ```c
   // コアレッシングされていないアクセスパターン
   int idx = threadIdx.x;
   float val = input[idx * stride];
   
   // コアレッシングされたアクセスパターン
   int idx = threadIdx.x;
   float val = input[idx];
  1. データレイアウト最適化: キャッシュミスを最小限に抑えるため、メモリ内のデータ構造を最適化して空間的な局所性を最大化する。これには、同じタイプのデータ要素をグループ化するstructure-of-arrays(SoA)レイアウトや、配列内の構造体を並べるarray-of-structures(AoS)レイアウトなどの手法が含まれます。Here is the Japanese translation for the provided markdown file, with comments translated but code left untranslated:

日本語訳:

  1. メモリレイアウト最適化: データ要素を同じ構造体に属するように保持することで、メモリアクセスのパターンを最適化するアプローチです。

    例:

    // 配列-構造体(AoS)レイアウト
struct Point {
    float x;
    float y;
    float z;
};
Point points[N];
 
// 構造体-配列(SoA)レイアウト
struct Points {
    float x[N];
    float y[N];
    float z[N];
};
Points points;

  2. キャッシュとプリフェッチ: メモリアクセスパターンの時間的、空間的なローカリティを活用することで、ハードウェアマネージドのキャッシュを効果的に活用できます。データタイリング (データを小さな塊に分割してキャッシュに収まるように変形する) やソフトウェアプリフェッチ (データが必要になる前にキャッシュにロードする) などの技術が使えます。

    例:

    // データタイリング
for (int i = 0; i < N; i += TILE_SIZE) {
    for (int j = 0; j < N; j += TILE_SIZE) {
        // キャッシュに収まるデータタイルを処理する
        for (int ii = i; ii < i + TILE_SIZE; ii++) {
            for (int jj = j; jj < j + TILE_SIZE; jj++) {
                // A[ii][jj]に対する計算を行う
            }
        }
    }
}

  3. メモリアクセススケジューリング: ROWバッファヒットを最大化し、DRAM プリチャージやアクティベーションのオーバーヘッドを最小化するようにメモリアクセスを並べ替えます。メモリコントローラーのハードウェア機構やアクセスパターン最適化、データレイアウト変換などのソフトウェア手法で実現できます。

  4. 圧縮: GPUコアとメモリ間のデータ転送量を削減するために、データ圧縮技術を適用します。これによりバンド幅の制約を緩和し、データ移動に関連するエネルギー消費を削減できます。

    例:

    • デルタ符号化: 連続する値の差分を保持する
    • ランレングス符号化: 繰り返し値を単一の値と個数で表現する
    • ハフマン符号化: 出現頻度の高い値により短いビット列を割り当てる

  5. メモリ仮想化:仮想化: GPUアプリケーションに統一された連続アドレス空間を提供するために、仮想メモリ技術を活用すること。これにより、メモリ管理をより柔軟に行えるようになり、デマンドページングなどの機能を実現できるため、メモリ使用量を削減し、システムの活用率を向上させることができます。

例:

  • CUDA における統一仮想アドレス (UVA): GPUスレッドがCPUメモリに単一のポインタでアクセスできるようにし、異種システムでのメモリ管理を簡素化します。

マルチチップモジュールGPU

GPUのパフォーマンスと電力要件が増大し続けるにつれ、従来の単一チップ設計では需要に追いつくことが難しくなってきています。そのため、複数のGPUチップを1つのパッケージに集積したマルチチップモジュール (MCM) 設計が、この問題に対する有望な解決策として登場してきています。

MCMGPUの設計には以下のような利点があります:

  1. 高メモリバンド幅: 複数のメモリスタックやチップを統合することで、単一チップ設計に比べて大幅に高いメモリバンド幅を提供できます。

  2. スケーラビリティの向上: MCM設計では、より多くのコンピューティングユニットとメモリコントローラを統合できるため、GPUのパフォーマンスをより高いレベルにスケールアップできます。

  3. 歩留まりと cost-effectiveness の向上: MCM設計では個々のチップが小さいため、大型の単一チップに比べて製造歩留まりが良く、よりコスト効率的になります。

一方で、MCMGPUの設計には以下のような新たな課題も生じます:

  1. チップ間通信: MCMパッケージ内の各チップ間の効率的な通信が、パフォーマンス上重要です。チップ間のデータ移動のオーバーヘッドを最小限に抑えるためには、高帯域幅かつ低レイテンシの相互接続が必要です。

  2. 電力供給と熱管理: MCM設計では、最適なパフォーマンスと信頼性を確保するために、慎重な電力供給と熱管理の戦略が必要となります。

  3. ソフトウェアサポート: MCMGPUでは、マルチチップアーキテクチャの利点を十分に活用するために、プログラミングモデルとランタイムシステムの変更が必要になる可能性があります。

この分野の研究は継続中です。以下がこのマークダウンファイルの日本語翻訳です。コードについては翻訳せず、コメントのみを翻訳しています。

MCM GPUの設計と最適化を探求するこの分野には、メモリシステムアーキテクチャ、インターコネクト設計、リソース管理などが含まれます。

例えば、Arunkumar et al. [2017]は、高帯域幅かつ低レイテンシのインターコネクトを使用してGPUチップを接続するMCM GPU設計を提案しています。また、MCM設計の高帯域幅とキャパシティを活用してパフォーマンスと energy 効率を向上させるメモリシステムアーキテクチャも提案しています。

別の例は、Milic et al. [2018]の研究で、リソース利用率を向上させ、チップ間通信オーバーヘッドを削減するためのリソース管理スキームを提案しています。このスキームは、ハードウェアとソフトウェアの組み合わせを使用して、アプリケーションのリソース使用状況と通信パターンを監視し、動的なリソース割り当てを行います。

結論

メモリシステムは現代のGPUアーキテクチャにとって重要な構成要素であり、その設計と最適化はシステム全体のパフォーマンスと効率に大きな影響を与えます。並列ワークロードの要求が増え続けるなか、研究者はGPUメモリシステムのパフォーマンス、スケーラビリティ、適応性を向上させるための幅広い手法を探求しています。

この分野の主な研究方向には、メモリアクセススケジューリングとインターコネクト設計、キャッシュの有効性、メモリリクエストの優先順位付けとキャッシュバイパス、ワーク間の異質性の活用、コーディネートされたキャッシュバイパス、アダプティブなキャッシュ管理、キャッシュの優先順位付け、仮想メモリページの配置、データ配置、マルチチップモジュールデザインなどがあります。

これらの手法を探求することで、研究者はパフォーマンスと energy 効率を維持しつつ、並列ワークロードの増大する需要に対応できるGPUメモリシステムを開発することを目指しています。GPUが機械学習、科学計算、データ分析などの新しい分野で進化し続けるにつれ、そのメモリシステムの設計と最適化がますます重要になっています。以下は日本語への翻訳です。コードに関しては、コメントのみを翻訳しています。

幹細胞は引き続き研究と革新の重要な分野であり続けるだろう。

Chapter 4 Gpu Memory System DesignChapter 6 Gpu Performance Metrics and Analysis