世界一のスパコン「神威 太湖之光」のプログラミングモデル

TaihuLightのプログラミング

TaihuLightのプログラミングモデルは、ノード間はMPIでノード内はOpenACC、あるいはAthreadと呼ぶライブラリで並列プログラムを書くというものである。MPIは超並列のスパコンのプログラミングではデファクトであり、ノード内はOpenACCではなくOpenMP4を使うとか、GPUであればOpenCLやCUDAを使うという手もあるが、TaihuLightのプログラミングモデルは標準的なものである。

なお、MPIは1つのMPEで1プロセスを動かしている。つまり、SW26010には4個のMPEがあるので、4つのMPIプロセスが動き、MPI的には4ノードになる。

OpenACCはGPUでも用いられているもので、CあるいはFortranのソースプログラムに並列化指示行を書き加えると、コンパイラが並列コードを生成してくれるというものである。

TaihuLightのOpenACCはデータ転送DMAの起動がMPEではなくCPEが行うなど標準的ではない点があり、また、高性能化のためのフレキシブルなDMAモードなどの機能を持っているので、それらに対応してカストマイズされていると考えられる。

Athreadは標準のPthreadに対応するものであるが、これもTaihuLightのハードウェアに合わせたカスタマイズが行われていると考えられる。このAthreadはSunway OpenACCを実装するための基本コンポーネントとして使われている。

Sunway OpenACCはOpenACC2.0規格に基づくコンパイラで、並列化指示行を加えたCあるいはFORTRANプログラムを、基本コンパイラで使えるCプログラムあるいはFOETRANソースプログラムに変換してくれる。このため、面倒なCPEへのデータ転送を一々プログラムの記述する必要がなくなり、並列プログラミングの労力を低減することができる。

このコンパイラはローレンスリバモア国立研究所(LLNL)で開発されたROSEコンパイラ基盤をベースに開発したものである。

次の図のように、標準のCプログラムに#pragma acc parallel loop copyin(B,C) copyout(A)という並列化指示行を追加してSWACCでコンパイルすると、MPEにはCPEでプロセスを起動するコードが作られ、CPEのコードにはメインメモリからCPEへのデータ転送、ループ、そして結果をメインメモリに転送するコードが作られる。

これを基本のコンパイラでコンパイルするとTaihuLight用のa.outが作られる。

ただし、普通のOpenACCでは、メインメモリからアクセラレータメモリへのデータ転送はホストCPU主導で行われるが、TaihuLightでは、CPEへのデータ転送がCPE主導で行われるという点が異なっている。

Sunway OpenACCでは、data copyデータ環境指示がアクセラレータの並列リージョンの中でもで使える、copy on parallel指示がスクラッチパッドメモリ(Local Data Memory)の間でのデータ移動や分配に使えるという拡張が行われている。また、pack/packin/packoutやswap/swapin/swapout句が新設されており、データ転送の効率を改善できるようになっている。

そして、データ移動と計算の並列実行をより良くコントロールするtilemask、entire、co_compute句などが追加されている。

AthreadはPOSIXのPthreadと似た機能を持つライブラリ群である。athreadライブラリを初期化するathread_init、スレッドを起動するathread_create、すべてのCPEを使って並列実行を行うathread_spawn、指定したスレッドの終了を待ち合わせるathread_waitなどの機能を持っている。

次の例は、MPEのコードでathread_initを呼び出して初期化を行い、次にathread_spawnでCPE側のスレッドを起動する。そして、athread_joinでspawnしたすべてのスレッドの終了を待ち合わせて、処理結果を印刷するコードである。CPE側でやっている計算は、配列の各要素の値を2倍するという処理である。