ヘテロジニアスな環境におけるソフトウェア開発
Agenda 今日の概要 ヘテロジニアスな環境の登場 ホモジニアスからヘテロジニアスへ ヘテロジニアスなアーキテクチャ GPU CUDA OpenACC, XeonPhi 自分のプログラムを理解するために デバッガ 共通の操作体験 TotalView 続きはブースで
より速く ホモジーニアスな並列 HPC 銀河生成 金融のリスク計算 車の衝突解析 製薬 大規模 複雑なアルゴリズム 高速 単一のCPUコアはクロック数が頭打ち ~ 3GHz ホモジーニアスな並列 マルチスレッド (OpenMP, pthreads) CPUクラスタ マルチプロセス (MPI) 高い汎用性 長年の実績とノウハウ メモリアクセス速度のボトルネック 高価なインターコネクト 複雑な非同期プログラミング The Free Lunch Is Over By Herb Sutter
専用の外付けデバイス より速く シンプルで高速な演算器 高い並列化効率 費用対効果 省電力 新しい言語拡張の理解 ハードウェアの理解 GPU アクセラレータ ヘテロジニアスな開発環境 MIC コプロセッサ
Top 500, Green 500, より速く
GPU と XeonPhi ヘテロジニアスな環境 GPU CUDA と OpenACC Device, SM, CUDA Core, Warp, Lane block, thread, global memory, shared, Xeon Phi offload ディレクティブ symmetric モード native モード 512bit のベクトル
ハードウェア ホストCPUのアクセラレータとして あるいはホストをGPU 演算の補助装置として 数千コアの並列 グリッド ( デバイス ) SM CUDAコア (SP) 複数 CPU x 複数 GPU 製品ライン GeForce, Quadro, Tesla, アーキテクチャ Kepler, Fermi, Maxwell, GPU アーキテクチャ
GPU アーキテクチャ ハードウェア階層 Grid SM (16 個 / Grid) (Fermi) CUDAコア (SP, 32 個 / SM) (Fermi) Warp( 縦糸, 32スレッド ) 制御単位 ベクトル Thread ( 撚り糸, 3D, レーン ) CUDAコアとregisterの組 軽量 メモリの種類 オフチップ Video memory オンチップ shared memory (64kB/SM) register constant cache, texture cache SM (Streaming Multiprocessor) Grid CUDA コア
CUDA NVIDIA CUDA GPUベンダーのNVIDIA 公式 ( 信頼 ベンダーロック ) 無償の開発環境 cufftやcublasなどの数値計算ライブラリ付属 OpenCV, MATLAB, IMSLなどツール経由でCUDA 利用することも ホストはWindows, Linux, Mac レイヤーが低く コードが複雑 CUDA 5.0 5.5 Dynamic Parallelism Cuda 6.0 Unified Memory
CUDA の論理階層 グリッド ( デバイス ) スレッドブロック (2D/3D) SMへの割付単位 スレッド (3D) CUDA アーキテクチャ Kernel 関数 ホスト側 ( 呼び出し ) MatMulKernel<<<dimGrid, dimblock>>>(mata, matb, matc); デバイス側 ( 実装 ) global void MatMulKernel(Matrix A, Matrix B, Matrix C) {... } 論理構造を明示的に指定
メモリ階層 ( 論理的 ) 開発者が明示的に指定 Memory Scope Locality Global Device External Shared Block Chip Local Thread Chip Constant Device Chip (cache) Texture Device Chip (cache) Register Thread Chip shared 16kB CUDA アーキテクチャ
CUDA プログラミング 言語 C 言語の拡張 コンパイラNVCC Python,Javaなどのバインディング ホストコードとデバイスコード 同一ファイル内に混在可能 NVCCが振り分けてくれる メモリ修飾子 device, constant, shared 関数修飾子 global, device, host
CUDA プログラミング 単純なプログラムの例 CUDA サンプル vectoradd 2つの配列の足し算 データを計算単位に分割する 組み込み変数 blockidx, threadidxごとに分ける 何番目のBlock (3D), 何番目のThread (3D) にいるのか?( 論理階層 ) 何番目のdevice, SM, Warp, Laneにいるのか?( 物理階層 ) kernel 関数の定義 ホストメモリ確保 ホストデータ初期化 デバイスメモリ確保 メモリ転送 kernel 関数呼び出し デバイスメモリ解放 ホストメモリ解放
OpenACC OpenACC の特徴 Fortran, C/C++ 有償 ディレクティブ指向 OpenMPとの類似 ホストコードのみ記述 通常のCPUコードとしても使える コンパイラが多くの仕事を担当 ループの検出 GPUの側のメモリ管理 CPUとGPUの間のデータ移動 Kernel 関数を作成するかどうか オープン規格 Cray, PGI, NVIDIAがサポート CUDAと組み合わせることができる OpenACC CUDA gang threadblock worker warp vector warp 内のスレッド OpenACC の論理階層
Xeon Phi Intel Xeon Phi コプロセッサ ホストCPU にPCI Expressで接続するボード MICアーキテクチャのx86 互換のコプロセッサ 汎用的なCPU 用に書かれたコードの移植性が高い コアごとに512bitのSIMD 処理 (16 命令 /clock, ベクトル長 ) 独立したOS(Linux) が動作し sshなどによるアクセス 柔軟な実行モデル Offloadモデル ディレクティブ 必要な部分だけコプロセッサに投げる シンメトリックモデル ホストとコプロセッサの間でMPIなどを使ってやりとりする Nativeモデル コプロセッサ上でのみ実行 既存のコードをそのまま使える
Xeon Phi offload のコード例 host offload void test08() { float pi = 0.0f; const int count = 10000; int i; OpenMP MIC t0 t1 t2 t3 MIC #pragma offload target (mic) #pragma omp parallel for num_threads(4) private(i) reduction(+:pi) for (i = 0; i < count; ++i) { float t = (float)((i + 0.5f) / count); pi += 4.0f / (1.0f + t * t); } pi /= count; } host
コードを理解する 様々なヘテロジニアス環境 ハードウェアやプログラミングモデルの特性を理解 メモリの制約を理解してエラーを回避する 複数言語 複数パラダイム クラスタ スレッド ヘテロジニアス vector 自分のコードがどう振る舞うか プログラムは意図したとおりではなく書かれたとおりに動く At Operation. 仕様書ではなくコードが全て 複数人による作業 引き継ぎ 移植 γνῶθι σεαυτόν ( 汝自身を知れ, Know thyself)
コードを理解する 自分のコードと仲良くなるために code reading, ペアプログラミング ベアプログラミング どの処理がデバイス上でどのように振舞っているか想像する デバッグはコーディングの2 倍大変 開発時間の制約 手ぶらで立ち向かうのは危険 効率的なデバッガ
TotalView TotalView 幅広いコンパイラ プラットフォーム C, C++, Fortran 77 & 90 Unix, Linux, MacOS X, ラップトップからスパコンまで 並列デバッグ マルチスレッド MPIデバッグ CUDA, Intel Xeon Phi, OpenACC メモリ デバッグ機能 : MemoryScape リバース デバッグ機能 : ReplayEngine パワフルで軽量 使いやすい GUI パッチ機能 breakpoint: Evaluation Point スクリプト用の CLI, バッチ デバッグ リモートデバッグ
TotalView 世界中の企業や研究機関 HPCwire 影響力のあるミドルウェア 20 (2014/6)
TotalView 様々な動作画面 OpenACC CUDA ( デバイスコード ) ホスト XeonPhi コプロセッサ 同じインターフェイス breakpoint, ステップ実行, 関数の呼出履歴 配列の表示,
CUDA のデバッグ画面 TotalView CUDA グリッドとブロックの次元 warp/ レーン warp/sm warp 数 GPU フォーカススレッドの論理座標 スタックトレースとインライン関数 パラメータ レジスタ 局所変数と共有変数 warp の PC
TotalView CUDA:2 つの座標をマッピング Grid, SM, Warp, Lane Grid, Block, Thread 物理座標 論理座標
Summary ヘテロジニアスな環境 開発は大変 開発者や研究者はGPU, Xeon-Phiなどの新しいデバイスになじまなければならない 同時に従来のOpenMPやMPIの手法も知っておく必要がある それぞれの環境ごとのプログラムの挙動の違いをイメージする これらを限られた時間の制約の中で行わなければならない コードを知るにはデバッガが便利! TotalViewは複数のホストとデバイス 幅広いアーキテクチャに対し同一の操作体験 Accelerating Great Code 結論 TotalView で定時に帰ろう!!
XeonPhi http://www.eetindia.co.in/art_8800700109_1800012_nt_af7fa4b9_2.htm
Xeon Phi The spectrum of models CPU 中心 マルチコアのホスト offload Intel Xeon Phi 中心 シンメトリック メニーコアのホスト 汎用的なシリアル および並列計算 ディレクティブで並列化 対等な MPI ネイティブ ホスト コプロセッサ Main( ) Foo( ) MPI_*() Main( ) Foo( ) MPI_*() Foo( ) Main( ) Foo( ) MPI_*() Main( ) Foo( ) MPI_*() Main() Foo( ) MPI_*() PCIe
How can TotalView help you? Effective Debugging requires the capability to control and examine specific instances of program execution in detail Threads and/or MPI Deadlocks and hangs Race conditions It provides Asynchronous thread control Powerful group mechanism Fortran and/or C++ Complex data structures Diving and recursive dive STL Collection Classes STLView Rich class hierarchies Powerful type-casting features Memory Analysis Leaks and Bounds Errors Automatic error detection tools Out of Memory Errors Analysis of heap memory usage by file, function and line Numerical errors Extensible data visualization Slicing and filtering of arrays Powerful expression system Conditional watchpoints TotalView provides an answer to the question : What is my program really doing?
TotalView NVIDIA CUDA CUDA 4.2, 5.0, 5.5, 6.0 対応 CUDAの統一メモリ (unified memory) CUDAの動的モードのプログラム 1つのセッション内でホストとデバイスのコードをデバッグ TeslaやFermiなどのハードウェア上のCUDA LinuxおよびGPUデバイス上のスレッドを可視化 デバイス ブロック スレッドメモリの階層構造を完全に可視化 デバイスのスレッドを論理座標とデバイス座標の両方で操作可能 CUDAの関数呼び出し ホストのピン止めされたメモリ領域 CUDAコンテクストを可視化 CUDAの関数をインラインでもスタック上でもハンドリング 使いやすい軽量なGUIと 自動化に適したCLI( コマンドライン ) 複数のNVIDIAデバイスを使うアプリケーション CUDAで高速化されたクラスタ上のMPIアプリケーション 統一された仮想アドレスとGPUDirect メモリエラーを検知してレポート CUDAの例外をハンドリング
CUDA のデバッグ画面 TotalView スレッド (x,y,z) ブロック (x,y,z) ボックスの中にある行番号をクリックして breakpoint を置く GPU フォーカス スレッド セレクタで CUDA スレッドのブロック (x,y) やスレッド (x,y,z) インデクスを変更する TotalView のスレッド ID 0 以上 : ホストスレッド 0 より小さい : GPU スレッド
TotalView CUDA の例 GPU フォーカススレッドの論理座標 CUDA グリッドとブロックの次元 レーンあたりの warp SM あたりのワープ ワープ数などなど スタックトレースとインライン関数 パラメータ レジスタ 局所変数と共有変数 warp の PC を指す矢印
TotalView GPU デバイスの情報を表示 デバイス情報 論理情報
TotalView OpenACC ホストCPU/GPUどちらでもステップ実行 コンパイラ Cray CCE 8.x OpenACC
TotalView Intel Xeon Phi シンメトリックモデルのデバッグができる初め てのデバッガ 付属のMemoryScapeでネイティブおよびシンメトリックモデルのメモリデバッグ ( オフロードは8.13では未対応 ) ホストとコプロセッサ両方のスレッドを完全に可視化 MPI プログラムの完全サポート オフロードコードによる異種混合アプリケーションのシンメトリックデバッグ Xeon Phi ネイティブ アプリケーションのリモートデバッグ Xeon および Xeon Phi 両方の非同期スレッド制御 マルチホスト マルチカード リバースデバッグは未対応
TotalView 1 つのデバッグセッションで同じ操作体験 ホスト コプロセッサ