大規模並列プロセッサのプログラミング：実践型アプローチ

📚 コンテンツ概要

このコースでは、CUDA C環境を用いたGPUコンピューティングおよび並列プログラミングの包括的な入門を提供します。GPUアーキテクチャ、データ並列性、スレッド管理、メモリ最適化、そして高度なパフォーマンスに関する考察について、MRI再構成や分子可視化といった実世界の事例を通じて解説します。

実践的で手を動かす形のガイドを通じて、高性能並列コンピューティングの技術をマスターしよう。

著者: David B. Kirk, Wen-mei W. Hwu

謝辞: Ian Buck, John Nickolls, NVIDIA DevTechチーム、Jensen Huang, David Luebke, Bill Bean, Simon Green, Mark Harris, Manju Hedge, Nadeem Mohammad, Brent Oster, Peter Shirley, Eric Young, Cyril Zeller.

🎯 学習目標

1. マルチコアCPUとマニーコアGPUの設計哲学および性能の発展軌跡の違いを明確に識別する。
2. 現代的なGPUアーキテクチャの主要な構成要素（ストリーミングマルチプロセッサ（SM）やメモリ構造など）を特定する。
3. アムダールの法則を適用し、理論的なスピードアップを計算し、直列処理のボトルネックの影響を把握する。
4. 固定機能パイプラインとプログラマブル統合プロセッサアレイのアーキテクチャ上の相違点を対比する。
5. 「GPGPU」が中間段階として果たした役割と初期のシェーダー・プログラミングモデルの制限を説明する。
6. 原子的操作、バリア同期、倍精度サポートなどのハードウェア機能が、スケーラブルな汎用コンピューティングへの移行を可能にした仕組みを分析する。
7. 行列乗算アルゴリズムにおけるデータ並列性を識別し、活用する。
8. デバイスメモリ管理（割り当て、ホストとデバイス間のデータ転送、解放）を実装する。
9. 適切なスレッドインデックスおよびグリッド/ブロック構成を使用して、CUDAカーネルを構築し起動する。
10. 多次元スレッド階層（グリッドとブロック）を設計し、複雑なデータ構造をGPUハードウェアにマッピングする。

🔹 レッスン1：並列コンピューティングとGPUアーキテクチャの基礎

概要: このレッスンでは、CPUとGPUの設計哲学の分岐によって引き起こされた、順次処理から並列処理への根本的な転換を学びます。学生は「マルチコア」と「マニーコア」の進化パターンを検討し、GPUが巨額のスループットを達成できるハードウェアアーキテクチャの理解を深め、アムダールの法則によるスピードアップの数学的制約を学びます。

学習成果:

• マルチコアCPUとマニーコアGPUの設計哲学および性能の発展軌跡の違いを識別する。
• ストリーミングマルチプロセッサ（SM）やメモリ構造を含む現代的GPUアーキテクチャの主要な構成要素を特定する。
• アムダールの法則を適用し、理論的なスピードアップを計算し、直列処理のボトルネックの影響を明らかにする。

🔹 レッスン2：GPUコンピューティングの進化と将来展望

概要: このレッスンでは、グラフィックス処理ユニット（GPU）のアーキテクチャの進化を追跡します。三角形レンダリング専用の固定機能ハードウェアとして始まった歴史から、現在の強力で統合された汎用並列プロセッサへと進化した過程を学びます。学生は硬直的なグラフィックスパイプラインからプログラマブルシェーダーへのシフト、GPGPU運動の登場、そして現在の科学技術シミュレーションを支えるスケーラブルなアーキテクチャの発展について探求します。

学習成果:

• 固定機能パイプラインとプログラマブル統合プロセッサアレイのアーキテクチャ上の相違点を対比する。
• 「GPGPU」が中間段階として果たした役割と初期のシェーダー・プログラミングモデルの制限を説明する。
• 原子的操作、バリア同期、倍精度サポートといったハードウェア機能が、スケーラブルな汎用コンピューティングへの移行を可能にした仕組みを分析する。

🔹 レッスン3：CUDAプログラム構造とメモリ管理

概要: このレッスンでは、CUDAプログラムの基本的なアーキテクチャについて学び、ホスト（CPU）とデバイス（GPU）での実行の違いを強調します。学生は行列演算におけるデータ並列性を識別し、CUDA APIを使って別々のメモリ空間を管理する方法を学び、単一プログラム複数データ（SPMD）スタイルを用いて、グリッド、ブロック、スレッドの階層構造により並列実行を組織する方法を習得します。

学習成果:

• 行列乗算アルゴリズムにおけるデータ並列性を識別し、活用する。
• 割り当て、ホストとデバイス間のデータ転送、解放を含むデバイスメモリ管理を実装する。
• 適切なスレッドインデックスおよびグリッド/ブロック構成を使用して、CUDAカーネルを構築し起動する。

🔹 レッスン4：高度なCUDAスレッドとスケジューリング

概要: このレッスンでは、CUDAにおけるスレッドの階層的組織について探求し、多次元インデックスが物理的なデータおよびハードウェアリソースにどのようにマッピングされるかに焦点を当てます。バリア同期のメカニズムと透明なスケーラビリティの詳細を説明し、高性能コンピューティングにおけるラテンシー耐性を実現するためのスレッド割り当てとワープベーススケジューリングのアーキテクチャ的原則をまとめます。

学習成果:

• 多次元スレッド階層（グリッドとブロック）を設計し、複雑なデータ構造をGPUハードウェアにマッピングする。
• ビルトインのCUDA変数（blockIdx、threadIdx、blockDim）を用いて正確なデータインデックスを実装する。
• バリア同期を適用し、異なるGPUアーキテクチャ間で透明なスケーラビリティを維持しながらデータ整合性を確保する。

🔹 レッスン5：メモリ最適化と共有メモリタイリング

概要: このレッスンでは、メモリ帯域幅とリソース制約が並列コンピューティングにおける主なボトルネックとなる理由を解説します。タイリング技術を用いたグローバルメモリトラフィックの削減方法を詳述し、同期バリア（__syncthreads()）の重要性と、レジスタと共有メモリの戦略的選択がパフォーマンス最適化に与える影響について説明します。

学習成果:

• レジスタと共有メモリの制限が、カーネル内の並列度（オキュパシー）に与える影響を分析する。
• タイリング技術によって達成されるグローバルメモリ帯域幅消費の削減量を定量的に評価する。
• 共有メモリアクセス時にデータ整合性を維持するために同期関数が必要不可欠であることを特定する。

🔹 レッスン6：パフォーマンス解析とSIMT実行

概要: このレッスンでは、CUDAカーネルの最適化に不可欠なアーキテクチャおよびアルゴリズム的考慮事項について探求します。単一命令複数スレッド（SIMT）ユニットとワープ分割という基本的な実行モデルから始め、メモリコアレス、タイルド行列乗算、ストリーミングマルチプロセッサ（SM）リソースの動的分割といった高度なパフォーマンスチューニング技術へと進みます。

学習成果:

• 多次元スレッドブロックがハードウェアの線形ワープ実行順序にどのようにマッピングされるかを分析する。
• 並列還元アルゴリズムにおける制御フローの分岐を評価し、最小化する。
• メモリコアレスとタイルドデータアクセスパターンを実装することで、グローバルメモリ帯域幅を最適化する。

🔹 レッスン7：浮動小数点演算と数値精度

概要: このレッスンでは、浮動小数点数の基本的なアーキテクチャについて学び、IEEE 754標準の要素である符号、過剰符号化指数、正規化仮数に焦点を当てます。学生はこれらのビットパターンが離散的な数直線上にどのようにマッピングされるかを探索し、この表現の制限が大規模な和集合のような複雑なアルゴリズムの精度に与える影響を理解します。

学習成果:

• 正規化表現と過剰符号化を用いて、ビットパターンから数値を計算する浮動小数点形式を分解する。
• 数直線上での表現可能な数の分布を視覚化し、指数と仮数に割り当てられるビット数の影響を説明する。
• ULPを用いて数値の不正確さを定量的に評価し、異なる丸めモードが誤差に与える貢献を特定する。

🔹 レッスン8：事例研究：MRI再構成の並列化

概要: このレッスンでは、GPU上で高級磁気共鳴画像（MRI）再構成を並列化する方法について学びます。非カートーシアン軌道向けの反復的再構成プロセスに注目し、F^H dカーネルの計算負荷を低減するために、ループ変形、定数メモリ管理、データレイアウトの再編成、およびハードウェア加速された三角関数関数の使用を最適化します。

学習成果:

• カートーシアンFFTベースの再構成から、非カートーシアンk-spaceデータ向けの反復的線形ソルバーに基づくアルゴリズムへの移行を理解する。
• ループ分裂とループ入れ替えを適用し、順次的なCコードを巨大なCUDAスレッドマッピングに適した構造に変換する。
• 定数メモリのチャンク化と配列型構造（AoS）データレイアウトを用いてメモリスループットを最適化する。

🔹 レッスン9：事例研究：分子可視化とマルチGPU実行

概要: このレッスンでは、GPUコンピューティングの実際の応用として、直接クーロン和（DCS）法を用いた電気的ポテンシャルマップの計算に焦点を当てます。学生は基本的なカーネル実装から始まり、命令アンロール、メモリコアレス、パディングを活用した高度に最適化されたバージョンまで進みます。

学習成果:

• CUDA定数メモリとグローバルメモリのラテントサイクリング隠蔽技術を用いて、直接クーロン和（DCS）カーネルを実装する。
• 命令アンロールと共通座標計算の再利用を通じて、カーネルのパフォーマンスを最適化する。
• メモリコアレスとパディング戦略を適用し、GPUグローバルメモリアクセスを最大帯域幅に合わせて整列させる。

🔹 レッスン10：計算的思考と並列アルゴリズムの選択

概要: このレッスンでは、順次的思考から並列問題解決への転換に焦点を当て、並列プログラミングの目的とアルゴリズムの戦略的選択について学びます。学生は問題を並列処理可能な単位に分解する方法、ドメイン科学とハードウェアアーキテクチャのギャップを埋めるための計算的思考を身につけ、アルゴリズムのパフォーマンスを評価する方法を習得します。

学習成果:

• 並列プログラミングの主な目的を特定し、アムダールの法則を用いて理論的なスピードアップを計算する。
• タスクレベルとデータレベルの分解の違いを区別し、原子中心（スキャッタ）とグリッド中心（ガザー）の戦略を適用する。
• メモリ帯域幅、計算複雑性、アーキテクチャ制約などの基準に基づいて、並列アルゴリズムを選定・評価する。

🔹 レッスン11：OpenCLプログラミングモデルの導入

概要: このレッスンでは、異種並列コンピューティングのフレームワークとしてのOpenCLを紹介し、そのデータ並列性モデルと階層的なハードウェア抽象化に焦点を当てます。学生は、OpenCLのNDRangeおよびメモリ構造をCUDAの同等物にマッピングする方法を学び、ランタイムでの動的コンパイルモデルを用いたホスト側のデバイス管理を習得します。

学習成果:

• OpenCLの並列性とメモリ階層を、CUDA固有のアーキテクチャ（例えばワークグループをブロックに、ローカルメモリを共有メモリにマッピング）にマッピングする。
• OpenCLカーネル関数を実装し、コンテキストとコマンドキューを使用してホスト側の実行環境を管理する。
• ソースコードからランタイムにカーネルをビルドする動的コンパイルワークフローを実行する。

🔹 レッスン12：現代のGPU機能と将来展望

概要: このレッスンでは、GPUのアーキテクチャ的および機能的な進化について探求し、高度なメモリ管理、強化されたカーネル実行機能、増加するコア性能への移行に焦点を当てます。学生は、ユニファイドデバイスメモリスペースやカーネルレベルの関数呼び出しといった機能が、GPUを汎用プロセッサへと進化させることの意義を検討します。

学習成果:

• メモリアーキテクチャの進化の意義と、64ビットユニファイドデバイスメモリスペースへの移行を説明する。
• 強化された原子操作とカーネルレベルの関数呼び出しが、複雑なデータ構造およびアルゴリズムの実装を可能にする仕組みを分析する。
• 同時カーネル実行、倍精度速度の向上、制御フロー効率の改善が、現代のGPU環境におけるパフォーマンスに与える影響を評価する。