OpenACCを活用するうえで重要な4つのステップ

OpenACCはディレクティブで並列化を指示する

OpenACCは、ソースプログラムにディレクティブと呼ぶ指示を追加することにより、コンパイラにCPUとGPUで動作するコードやデータのコピーを行うコードを生成する。図4に示すように、赤字で書いた#pragma acc parallel loopという行を挿入すると、OpenACCをサポートするコンパイラは、その行の次のfor文をGPUで並列化して実行する実行するコードとして生成し、GPUメモリに必要な領域を確保し、必要なデータをコピーするコードと結果をCPUメモリにコピーするコードを付け加えてくれる。

結果として、#pragma以前の文とfor文の後の文はCPUで実行され、for文の中の文はGPUで実行される。つまり、プログラマは面倒なデータの移動を記述しないで済むし、GPUで動作する部分をCUDAで別途、カーネルプログラムとして書く必要も無くなる。

このプログラムは#pragmaの行を無視してコンパイルすれば、CPUだけで走るコードを生成でき、CPU用とCPU+GPU用に別のソースコードを作ってメンテする手間も掛からない。

こういうと良いことずくめであるが、一般にCUDAでチューニングされたプログラムの性能には及ばないことが多いようである。

OpenACC化の最初のステップは並列性の発見

ディレクティブを使うプログラミングは、図5に示す、並列性の発見、並列性の記述、データ局所性の記述、最適化の4ステップで行うのが良い。また、必要に応じて、この4ステップを繰り返すことによって、さらに、性能を上げることができる。

以下では、ヤコビ法のプログラムを例にとり、OpenACCプログラミングの基礎を説明していく。

この計算は、図6に示すように周囲の4点の値の平均を中央の点の次の値とするという計算を、全部の点について行う。そして、各点の次の値が求まったら、それを基にして、次の値を計算するという処理を繰り返し、値が変わらなくなると収束で、それが解ということになる。

この処理を行うコードは図7のようになる。最初のWhile文は収束するまでの繰り返しを行う。次の2重のfor文は各点の次の値(Anew)を周囲の4点(A)の値の平均値として計算するループである。また、このループの中では元の値と次の値の差の絶対値を計算し、これをerrとする。errは全ての点のerrの最大のものを求めるため、以前の最大値とのMaxを計算している。

そして、すべての点の次の値(Anew)が求まると、それを元のAにコピーして、errがtolの値より小さくなり収束するまで、この処理を繰り返す。

このプログラムから、並列に実行できる部分を見つけるのが、最初のステップである。

最初のWhile文は前回のerrの値に依存しているので、whileの繰り返しは並列に実行することはできない。一方、その次の2重のforループは、i、jともにどのような順序で実行しても同じ結果が得られるので、並列化できるループである。また、そのあとのAnewからAにコピーする2重ループもどのような順序で実行してもよく、この部分も並列に実行できるループである。

OpenACC化の第2ステップは、並列性の記述

そして、この並列性をコンパイラに教えるために、図9に示すように、並列性のあるfor文の前に#pragma acc parallel loopを追加する。

最初のループではreduction(max:err)という指定がついている。この2重ループはGPUで実行するスレッドに分解されて並列に実行されるのであるが、reductionは、全スレッドのerrのmaxを最終結果として出力すべきことを指示している。後のAnewのコピーのループでは全スレッドの最大値を求めるというようなreductionがないので、最初のループのような指定は必要ない。

この#pragam acc parallel loopと書く記述は、OpenMPの記述と近く、OpenMPを使ってきたプログラマにはなじみやすい記述である。しかし、OpenACCには、図10に示す#pragma acc kernelsという書き方もある。図10のように書くと、コンパイラが2つのforループを認識して、2つのカーネルを生成し、データの受け渡しのコードも生成してくれる。OpenMPを使っていなかった人には、こちらの方が使いやすいかもしれない。

図11はヤコビ法の例をkernels指定で書いたもので、図9とまったく同じ結果になるとのことである。