第7回:今あるプログラムを楽に速くするためには

今回で 拡散方程式のプログラム 例は終わりです。

ここまで学んだこと(+α)のまとめとして、何をすればどのくらい速くなるのかを見ていきます。

最適化の効果を確認しよう

ここまでで解説した指示文(+α)を使って最適化した場合、どのくらい速くなるのか確認してみましょう。

図1のグラフは、東京大学のReedbushスパコン を使い、CPUとGPUの性能比較をしたものです。

CPUはIntel Xeon E5-2695v4 (Broadwell-EP) を1ソケット(18 コア)、GPUはNVIDIA Tesla P100を1基使っています。
​いずれも2016年の導入当時最新のもので、少し世代の古いものですが、おおよその傾向は最新のものでも大きくは変わりません。

図1:CPUとGPUの性能比較
図1:CPUとGPUの性能比較

図1の評価では、CPU向けの並列化はOpenMP、GPU向けの並列化はOpenACCまたはCUDAを使っています。

コンパイルにはそれぞれ、Intel compiler (icc –O3 –qopenmp)、PGI compiler (pgcc –O3 –acc –ta=tesla,cc60)、nvcc compiler (nvcc –O3 –arch=sm_60)を使っています。

グラフの縦軸はFlopsと呼ばれる指標で、一秒間に何回浮動小数点演算できたかを表しています。
つまり、数値が高い程、性能が良いということです。

左から2本のグラフが、CPUでの結果(1コアの場合と全18コアを使った場合)を表しています。
OpenMP版の実装では、#pragma omp parallel for の一行を挿入しただけですが、18コアを使うことによりおよそ14倍程度高速化しています。

一方で「GPU①:kernels, loop指示文で並列化」の実行では、初めからGPUの全コアを利用しています。​それにもかかわらず性能はCPUの1コアを下回っています。

これはここまでに説明してきた通り、CPUとGPUの間の冗長なデータ転送が発生してしまうためです(図2)。

図2:GPU①の実装
図2:GPU①の実装

これはGPU②(図3)のように、CPU-GPU間のデータ転送を最小化することで高速化できます。

​この時点で、CPUの18コア実装と比較して、13倍程度の性能を得られました!

図3:GPU②の実装
図3:GPU②の実装

ここまでは、今まで解説してきたことのおさらいです。

ここからは、OpenACC実装の性能をよりCUDA版に近づけるための+αについて、少しだけ解説します。

GPU③(図4)は、その+αとして、gang, vectorによるGPUのスレッド割り当ての調整、asyncによる、CPU側のコストの隠蔽を行っています。

Gang, vectorによるスレッド割り当ての調整については、第五回の記事で少し触れました。
その時にも述べた通り、これを適切に設定することは難しいです。

図4:GPU③の実装
図4:GPU③の実装

解きたい問題のサイズによっても変わりますので、基本的にはコンパイラの自動的設定(gang, vectorを書かない場合には自動で設定される)に頼ってよいと思います。

ただ、一般的に最も内側のループのvector値は32の倍数が良いです。
コンパイラのメッセージを見て、最も内側のループのvector値が32の倍数になっていないときは、利用を考えてみてください。

さてもう一方の asyncですが、これは非同期処理を行うための指示子です(図5)。
GPUのコアなどが余っているときに、別のカーネル同士を同時に実行するための指示子ですが、CPUがGPUのカーネル終了を待たなくなるという効果もあります。

実は、CPUがGPUを呼び出すときにかかる時間は馬鹿にならず、問題サイズが小さいときには特に性能に影響します。図4の最適化でasyncを使っているのは、このCPUがGPUカーネルを呼び出す時間を隠蔽するためでした。

今回の問題は1283サイズですが、1割程度性能に影響しています。なお、CUDAの場合デフォルトでもCPUはGPUの実行終了を待たないので、その分が OpenACC との差となっています。

図5:asyncによる非同期実行
図5:asyncによる非同期実行のイメージ

今回は、OpenACCにおける最適化がどの程度の効果を示すのか、OpenMPによるCPU実行や、CUDA実装との比較により示しました。
拡散方程式のように、OpenACC で簡単に並列化できるプログラムであれば、CUDAと遜色ない性能を発揮できることがわかっていただけたかと思います。

なお、今回は指示文の変更による最適化のみを行いましたが、OpenMP版、OpenACC版、CUDA版いずれも、プログラム本体を最適化することで、さらなる性能向上が見込めます。​

1ヵ月間有効のスパコンお試しアカウント

東京大学情報基盤センターでは、教育の一環として、制限はあるものの一ヵ月の間有効なスパコンアカウントを提供しています。

現在3つのスパコンが運用されていますが、そのうちReedbushと呼ばれるスパコンには、一世代前のものではありますがGPUが搭載されていて、OpenACCを使える環境も整っています。

自分でどんどん自習したい場合は、ご利用を考えてみてください。

トライアルアカウント申し込みページ
https://www.cc.u-tokyo.ac.jp/guide/trial/free_trial.php

< 過去の講習会の資料やプログラム公開中 >

東大センターが行った過去のOpenACCに関する講習会の資料やプログラムも公開されていますので、自習する場合にはぜひご利用ください。オンライン講習会 定期開催中!

講習会ページ
https://www.cc.u-tokyo.ac.jp/events/lectures/

講習会で用いているプログラム
https://www.dropbox.com/s/z4fmc4ibdggdi0y/openacc_samples.tar.gz?dl=0​