データ転送とGPU処理を並列に実行させられるAsync指定

データリージョンとupdate指定などをうまく使うことによって、無駄な転送を省き、プログラムの性能を上げることができるが、必要なデータ転送だけは残ってしまう。しかし、このデータ転送をGPU処理と並列に実行することができれば、転送時間はGPU処理に隠れて、実質的に転送時間ゼロとすることも可能となる。

このように、CPUの動作とGPUの動作を並行(非同期に)実行させるために、OpenACCにはAsyncという指定が設けられている。このAsyncは、ほとんどのOpenACCディレクティブにつけることができる。ディレクティブでparallel loopなどの指定があるとGPUカーネルが起動されるが、Asyncが付いていると、CPUはGPUカーネルの終了を待たず、CPUプログラムで指定された次の仕事を実行して行く。

Asyncが付いたディレクティブは、その終了を待たず、CPUは処理を続けていくが、また、Asyncのついたディレクティブが出てくると、その処理は、キューに入れられ、前のasyncの処理が終わったら実行されることになる。

このキューは複数作ることができ、async( )の引数はキューの識別子で、同じ識別子のasyncがついたディレクティブは、同じキューに入れられる。1つのasyncキューに入れられた処理は順に実行されるが、他の識別子のasyncキューの動作との関係は非同期になる。

Asyncの識別子は、単純に整数で番号を付けても良いし、使いかたによってはループのインデックスなどを識別子として使ってもよい。

図9 CPUとは非同期の動作を指定するAsync。識別子ごとにキューが作られ、1つのキューの中の動作は順番に実行される

そして、Asyncで非同期に実行しているGPUの処理の終了を待ち合わせるのが、図10に示したWaitディレクティブである。

#pragma acc waitと書くと、すべてのasyncキューの処理の終了を待ち合わせる。#pragma wait(識別子)のように、識別子を付けると、指定された識別子のasyncキューの処理の終了だけを待ち合わせ、他の識別子のasyncキューの処理の終了はチェックしない。

#pragma wait(識別子1)　async(識別子2)のようにWaitディレクティブにもasyncを付けることができる。このように書くと、識別子2のasyncキューは、識別子1のasyncキューのすべての処理が終わるまで待ち合わせるという意味になる。つまり、識別子1のキューに識別子2のキューを繋いだ処理となる。

図10 #pragma acc waitと書くと、すべての非同期処理の終了を待つ。Waitに識別子を付けると、そのキューの処理だけの終了を待つ

図11の例は、plane[p]と書かれた複数のプレーンに対して順に何らかの処理を行うもので、各プレーンは完全に独立で、CPUとGPUのメモリ間でデータの転送が必要であるとする。

図11 複数の独立のplane[p]の処理をオフロードして実行させる例

#pragma acc update device(plane[p])は、デバイスメモリのplane[p]の部分を、CPUからコピーして更新するというディレクティブである。そして、#pragma acc parallel loopで、plane[p]のデータを加工する処理をGPUに並列実行させている。そして、GPUの処理が終了すると、#pragma acc update self(plane[p])で、GPUメモリからCPU(self)のメモリにデータをコピーしてアップデートしている。

これを図12のように、3つのディレクティブにasync(p)を付けると、update device( )によるGPUへのデータ転送が起動され、CPUは次のparallel loopの実行に入る。しかし、このディレクティブにもasync(p)がついており、識別子が同じpであるので同じキューに入れられ、GPUでのloopの実行はデータ転送が終わってからとなる。

そして、CPUは、次のCPUへのデータ転送のディレクティブの処理にかかるが、これもasync(p)がついているので同じキューに入れられ、その処理の開始はGPUでのloopの処理が終わってからということになる。

さらに、CPUはforループを回って、識別子をp+1として、update device( )を実行しようとする、この処理はasyncの識別子が異なるので、識別子pのキューの処理の終了を待たずに開始することができる。

なお、すべてのプレーンの処理が終わったことを待ち合わせるために最後に#pragma acc waitディレクティブが置かれている。

図12 3つのディレクティブにasync(p)を付けて非同期に処理を行わせる

結果として、図11のコードでは、Host to Deviceのデータ転送、処理、Device to Hostのデータ転送が順に行われ、仮にすべての処理の実行時間が同じとすると、図13の上側に示すように、2つのプレーンを処理するのに、全部で6単位の時間がかかる。

一方、図12のコードでは、識別子がpの処理は順番に行われるが、次の識別子p+1の処理は識別子pの処理が終わるのを待つ必要はない。このため図13の下の図のように、パイプライン的に処理を行うことができ、4単位の時間で2プレーンを処理することができる。実際にはこの3つの処理の長さは同じではないが、並列処理の基本的な考え方はあてはまる。

図13 図11のコードでは上の図のように順次実行となるが、図12のようにasyncを使うと、下の図のようにパイプライン処理ができる

大きなデータを複数のGPUで分担する処理を行う場合、ヤコビ法などでは、それぞれのGPUが分担する部分に隣接する外側の部分のデータも持つ必要がある。この外側の部分は、隣接する部分を分担するGPUのデータであり、この部分のデータをGPU間で転送する必要がある。

このような場合には、図12のコードのように転送と処理をパイプライン実行すれば、転送時間は処理時間に隠れてしまい、図11のコードのように順次処理する場合と比較して、実行時間を短縮することができる。

また、GPUにカーネルを転送し、実行を行えるようにするためには時間がかかるが、次のカーネルの処理をキューに詰め込んでおくと、この作業を事前にやっておくことができる。このため、パイプライン処理ができない場合でも、GPUが次の仕事を次の仕事を開始するまでの所要時間が短くなり、小さなカーネルを次々と実行する場合には有利である。