第6章: GPU パフォーマンスメトリクスと分析

GPUアプリケーションのパフォーマンスを分析し最適化することは、GPUハードウェアリソースの高効率と高利用率を達成するために重要です。この章では、主要なGPUパフォーマンスメトリクス、プロファイリングとオプティマイゼーションツール、パフォーマンスボトルネックの特定手法、GPUパフォーマンス向上のための戦略について探っていきます。

スループット、レイテンシ、メモリ帯域幅

GPUパフォーマンスを評価する3つの基本的なメトリクスは、スループット、レイテンシ、メモリ帯域幅です。これらのメトリクスとその意味を理解することは、GPUアプリケーションの分析と最適化に不可欠です。

スループット

スループットとは、GPUが一定時間内に完了できる演算や処理の数を指します。一般的には1秒あたりの浮動小数点演算数(FLOPS)や1秒あたりの命令数(IPS)で測定されます。GPUは並列性を活用し、多数のスレッドを同時に実行することで高いスループットを実現するよう設計されています。

GPUの理論ピークスループットは以下の式で計算できます:

ピークスループット(FLOPS) = CUDA コア数 × クロック周波数 × 1 CUDA コアあたりのサイクルFLOPS

例えば、NVIDIA GeForce RTX 2080 Tiには4352個のCUDAコアがあり、ベースクロック周波数は1350 MHz、各CUDAコアは1サイクルあたり2つの浮動小数点演算(FMA - Fused Multiply-Add)を実行できます。したがって、その理論ピークスループットは:

ピークスループット(FLOPS) = 4352 × 1350 MHz × 2 = 11.75 TFLOPS

ですが、実際にこの理論ピークスループットを達成するのは、メモリアクセスパターン、ブランチ分岐、リソース制約などの要因により難しいです。

レイテンシ

レイテンシとは、単一の演算や処理が完了するまでの時間を指します。GPUの文脈では、メモリアクセス操作に関連付けられることが多いです。GPUには階層的なメモリシステムがあり、メモリ階層の異なるレベルからデータにアクセスすると、以下は、提供されたマークダウンファイルの日本語翻訳です。コードについては、コメントのみ翻訳しています。

GPU における各メモリレベルの典型的なレイテンシは以下の通りです:

レジスタ: 0-1 サイクル
シェアードメモリ: 1-2 サイクル
L1 キャッシュ: 20-30 サイクル
L2 キャッシュ: 200-300 サイクル
グローバルメモリ (DRAM): 400-800 サイクル

レイテンシは、演算間の依存関係や、メモリからのデータ取得待ちが発生する場合に、GPU のパフォーマンスに大きな影響を及ぼします。レイテンシ隠蔽、プリフェッチ、キャッシングなどの手法を使うことで、レイテンシの影響を軽減することができます。

メモリ帯域幅

メモリ帯域幅とは、GPU とメモリサブシステム間のデータ転送レートを表します。一般的には、バイト毎秒 (B/s) やギガバイト毎秒 (GB/s) で表されます。GPU は、グラフィックスやコンピューティングワークロードの高帯域幅要求に対応するため、GDDR6 や HBM2 などの高帯域幅メモリインターフェースを備えています。

GPU の理論ピークメモリ帯域幅は、以下の式で計算できます:

ピーク メモリ帯域幅 (GB/s) = メモリクロック周波数 × メモリバス幅 ÷ 8

例えば、NVIDIA GeForce RTX 2080 Ti GPU のメモリクロック周波数は 7000 MHz (実効値)、メモリバス幅は 352 ビットです。したがって、その理論ピークメモリ帯域幅は:

ピーク メモリ帯域幅 (GB/s) = 7000 MHz × 352 ビット ÷ 8 = 616 GB/s

多くの GPU アプリケーションはメモリ帯域幅に制限されるため、メモリ帯域幅は GPU のパフォーマンスに非常に重要です。メモリアクセスパターンの最適化、データ転送の最小化、メモリ階層の活用などによって、メモリ帯域幅の活用を改善することができます。

プロファイリングとパフォーマンス最適化ツール

プロファイリングとパフォーマンス最適化ツールは、GPU アプリケーションの動作を分析し、パフォーマンスボトルネックを特定し、最適化の指針を得るために不可欠です。これらのツールは、カーネル実行時間、メモリアクセス、スレッド実行状況など、GPU のパフォーマンスに関する様々な側面を提供します。以下は、提供されたマークダウンファイルの日本語翻訳です。コードについては、コメントのみ翻訳しています。

GPUのプロファイリングとパフォーマンス最適化のための人気のツールには以下のようなものがあります:

NVIDIA Visual Profiler (nvvp): GPUアプリケーションのパフォーマンスを包括的に表示するグラフィカルなプロファイリングツールです。開発者はカーネル実行、メモリ転送、APIコールを分析し、最適化のためのアドバイスを得ることができます。
NVIDIA Nsight: GPUアプリケーションのプロファイリングとデバッグ機能を含む統合開発環境(IDE)です。CUDA、OpenCL、OpenACCなどさまざまなプログラミング言語とフレームワークをサポートしています。
NVIDIA Nsight Compute: GPUカーネルのパフォーマンス分析に焦点を当てた独立したプロファイリングツールです。命令スループット、メモリ効率、占有率などの詳細なパフォーマンスメトリクスを提供し、ソースコードレベルのパフォーマンスボトルネックを特定するのに役立ちます。
AMD Radeon GPU Profiler (RGP): DirectX、Vulkan、OpenCLアプリケーションのパフォーマンスデータをキャプチャーし、可視化するAMD GPUのプロファイリングツールです。GPUの利用率、メモリ使用量、パイプラインのストールなどの洞察を提供します。
AMD Radeon GPU Analyzer (RGA): GPUシェーダーコードを静的に分析し、パフォーマンス予測、リソース使用量、最適化の提案を行うツールです。

これらのツールは通常、GPUアプリケーションコードにインストルメンテーションを行い、実行中にパフォーマンスデータを収集し、分析のためにユーザーフレンドリーな形式で表示します。タイムラインビュー、パフォーマンスカウンター、ソースコードとの相関関係を提供し、開発者がパフォーマンス問題を特定し、コードを最適化するのに役立ちます。

例: NVIDIA Visual Profiler (nvvp)を使ってCUDAアプリケーションをプロファイリングする

プロファイリングを有効にしてCUDAアプリケーションをビルドする:
```
nvcc -o myapp myapp.cu -lineinfo
```
プロファイリングを行いながらアプリケーションを実行する:
```
nvprof ./myapp
```
Visual Profilerを開く:
```
nvvp
```
生成されたプロファイリングデータをインポートする5. タイムラインビュー、カーネルパフォーマンス、メモリ転送、APIコールを分析します。
プロファイラの推奨事項に基づいて、パフォーマンスボトルネックを特定し、コードを最適化します。

パフォーマンスボトルネックの特定

GPUアプリケーションを最適化するためには、パフォーマンスボトルネックを特定することが重要です。パフォーマンスボトルネックは、非効率なメモリアクセスパターン、低いオキュパンシー、ブランチダイバージェンス、リソース制約など、さまざまな要因から生じる可能性があります。パフォーマンスボトルネックを特定するための一般的な手法には以下のようなものがあります:

プロファイリング: プロファイリングツールを使用してカーネル実行時間、メモリ転送時間、APIオーバーヘッドを測定することで、アプリケーションのどの部分が最も時間とリソースを消費しているかを特定できます。
オキュパンシーの分析: オキュパンシーとは、アクティブなウォープの数と GPU がサポートできる最大ウォープ数の比率を指します。低いオキュパンシーは GPU リソースの過小利用を示唆し、ブロックやグリッドの次元の最適化や、レジスタやシェアードメモリの使用量の削減が必要かもしれません。
メモリアクセスパターンの検討: 非コアレスメモリアクセスや、グローバルメモリへの頻繁なアクセスなど、非効率なメモリアクセスパターンは GPU パフォーマンスに大きな影響を及ぼします。プロファイリングツールを使ってメモリアクセスパターンを分析し、シェアードメモリの使用や、データローカリティの改善などの最適化の機会を特定できます。
ブランチダイバージェンスの調査: ブランチダイバージェンスは、ウォーピ内のスレッドが条件文によって異なる実行パスを取ることで発生します。これにより直列化が生じ、パフォーマンスが低下します。ブランチダイバージェンスを特定し、最小限に抑えることでGPUパフォーマンスを改善できます。
リソース利用状況の監視: GPUには、レジスタ、シェアードメモリ、スレッドブロックなど、限られたリソースがあります。プロファイリングツールを使ってリソース利用状況を監視し、リソースボトルネックを特定することで、レジスタ使用量の削減などの最適化につなげることができます。以下は、提供されたマークダウンファイルの日本語翻訳です。コードについては、コメントのみ翻訳しています。

共有メモリに収まるようにデータを分割または分割する。

例: NVIDIA Nsight Computeを使用したメモリアクセスボトルネックの特定

Nsight Computeを使ってCUDAアプリケーションをプロファイリングする:
```
ncu -o profile.ncu-rep ./myapp
```
生成されたプロファイルレポートをNsight Computeで開く。
"Memory Workload Analysis"セクションを分析して、非コアレスアクセスや大量のグローバルメモリ使用など、非効率なメモリアクセスパターンを特定する。
Nsight Computeが提供する洞察に基づいて、共有メモリの使用や data localityの改善など、メモリアクセスパターンを最適化する。

GPUパフォーマンス向上のための戦略

パフォーマンスボトルネックが特定されたら、さまざまな最適化戦略を採用してGPUパフォーマンスを向上させることができます。一般的な最適化戦略には以下のようなものがあります:

並列性の最大化: アプリケーションを十分な数の並列タスクに分解して、GPUリソースを最大限に活用する。これには、ブロックやグリッドの次元の調整、ストリームを使った並行実行、タスク並列性の活用などが含まれる。
メモリアクセスパターンの最適化: グローバルメモリアクセスの最小化、頻繁にアクセスされるデータの共有メモリの使用、コアレスメモリアクセスの確保など、メモリアクセスの効率を高める。メモリタイリング、データレイアウトの変換、キャッシングなどの手法が役立つ。
ブランチダイバージェンスの削減: ワープ内の発散ブランチを避けるようにコードを再構成する。ブランチ予測、データ依存ブランチ、ワープレベルプログラミングなどの手法でブランチダイバージェンスの影響を軽減する。
メモリ階層の活用: レジスタや共有メモリを使って頻繁にアクセスされるデータを最大限活用する。テクスチャメモリや定数メモリを、空間的局所性や均一アクセスを持つ読み取り専用データに活用する。
計算とメモリアクセスの重複: カーネル実行とメモリ転送の重複、ストリームを使った並行実行などにより、計算とメモリアクセスの重複を最大化する。以下は、提供されたマークダウンファイルの日本語翻訳です。コードの部分は翻訳せず、コメントのみ翻訳しています。
メモリ転送の隠蔽: CUDA ストリームや OpenCL コマンドキューを使用して、計算とメモリ転送を重複させることで、メモリ転送の待ち時間を隠蔽します。これにより、ホストとデバイスメモリ間のデータ転送中に、GPUが計算を実行できるようになります。
カーネルの起動パラメータのチューニング: ブロックサイズとグリッドサイズを変更して、各カーネルの最適な構成を見つけます。最適な起動パラメータは、スレッドあたりのレジスタ使用量、共有メモリの使用量、GPUアーキテクチャの特性などの要因に依存します。
ホストとデバイス間のデータ転送の最小化: できるだけGPU上で計算を行うことで、ホスト(CPU)とデバイス(GPU)間のデータ転送量を減らします。小さな転送をまとめて大きな転送にすることで、各転送のオーバーヘッドを軽減します。
非同期操作の活用: 非同期メモリコピーやカーネル起動などの非同期操作を活用して、計算と通信の重複を図ります。これにより、CPUは他のタスクを実行できるようになり、アプリケーション全体のパフォーマンスが向上します。

例: CUDA における共有メモリを使用したメモリアクセスパターンの最適化

非効率なグローバルメモリアクセスの元のコード:

__global__ void myKernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

共有メモリを使用して最適化したコード:

__global__ void myKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int localIdx = threadIdx.x;
 
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        dat
```a[tid] = result;
    }
}

最適化されたコードでは、入力データがまず共有メモリにロードされます。共有メモリはグローバルメモリに比べてアクセス遅延が非常に低いです。その後、共有メモリを使って計算が行われ、グローバルメモリへのアクセス回数が減らされ、パフォーマンスが向上します。

結論

GPUのパフォーマンスを分析し最適化することは、効率的で高性能なGPUアプリケーションを開発するために不可欠です。スループット、レイテンシ、メモリ帯域幅などの主要なパフォーマンス指標を理解することで、開発者は自身のコードの最適化について適切な判断ができるようになります。

プロファイリングとパフォーマンス最適化ツールは、パフォーマンスボトルネックを特定し最適化の方向性を示すのに不可欠な役割を果たします。これらのツールはカーネル実行、メモリアクセスパターン、占有率、リソース利用状況などについての貴重な洞察を提供し、開発者が最も重要な領域に最適化の努力を集中できるようにします。

GPUパフォーマンスを最適化するための一般的な戦略には以下のようなものがあります。Markdownフォーマットで続けます:

ブランチ発散の削減: ワープ/ウェーブフロント内の発散的な制御フローは直列化とSIMD効率の低下につながる可能性があります。アルゴリズムは可能な限りブランチ発散を最小限に抑えるように構造化されるべきです。ブランチ予測、データ依存ブランチ、ワープレベルプログラミングなどの手法がブランチ発散の影響を軽減するのに役立ちます。
メモリ階層の活用: レジスタと共有メモリを頻繁にアクセスされるデータに最大限活用することで、GPUメモリ階層を効果的に活用できます。テクスチャメモリと定数メモリは、空間的な局所性を持つか、スレッド間で一様にアクセスされる読み取り専用データに使用します。
計算とメモリ転送の重複: CUDAストリームやOpenCLコマンドキューを使ってメモリ転送と計算を重複させることで、メモリ転送レイテンシを隠蔽できます。これにより、以下は、提供されたマークダウンファイルの日本語翻訳です。コードの部分は翻訳せず、コメントのみを翻訳しています。
カーネルの起動パラメータのチューニング: ブロックサイズとグリッドサイズを変更して、各カーネルの最適な構成を見つけてください。最適な起動パラメータは、スレッドあたりのレジスタ使用量、共有メモリの使用量、GPUアーキテクチャの特性などの要因に依存します。
ホストとデバイス間のデータ転送の最小化: GPUで可能な限り多くの計算を行うことで、ホスト(CPU)とデバイス(GPU)間のデータ転送量を減らしてください。小さな転送をまとめて大きな転送にすることで、各転送のオーバーヘッドを軽減できます。
非同期操作の活用: 非同期メモリコピーやカーネル起動などの非同期操作を活用して、計算と通信を重複させてください。これにより、CPUは他のタスクを実行できるようになり、アプリケーション全体のパフォーマンスが向上します。

例: CUDAの共有メモリを使ったメモリアクセスパターンの最適化

非効率なグローバルメモリアクセスの元のコード:

__global__ void myKernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

共有メモリを使った最適化されたコード:

__global__ void myKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int localIdx = threadIdx.x;
 
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        data[tid] = result;
    }
}

最適化されたコードでは、入力データを最初に共有メモリにロードしています。共有メモリはグローバルメモリに比べてアクセスレイテンシが低いため、以下は、提供された Markdown ファイルの日本語翻訳です。コードについては、コメントのみを翻訳しています。

グローバルメモリ

計算は共有メモリを使って行われ、グローバルメモリへのアクセス数を減らし、パフォーマンスを向上させます。

// グローバルメモリからデータをロードする
data = global_memory[index];
 
// 共有メモリにデータをコピーする
shared_memory[thread_id] = data;
 
// 共有メモリを使って計算を行う
result = compute(shared_memory);
 
// 結果をグローバルメモリに書き込む
global_memory[index] = result;

Chapter 5 Gpu Memory System Design Chapter 7 Streaming Multiprocessor Design