プログラム階層とは独立の非構造データリージョン

普通のデータリージョンは#pragma acc dataに続く{ }の中がその範囲となるが、このようなプログラム階層とは関係なくデータリージョンを作る機能がUnstructured Data Regionで、データリージョンの開始の指定と終了の指定を個別に書くことができる。

図4の左側の図は普通のdata regionで、#pragma acc data copy( )の次の{ }の中が1つのデータリージョンになり、開始時にcopyin、終了時にcopyoutが行われる。これを右側の図のように、非構造データリージョンでは、開始時には#pragma acc enter data copyin( ) 、終了時には#pragma acc exit data copyout( )と書く。

図4 Unstructured Data Regionはデータリージョンの開始と終了を独立に指定できる

この例では得をすることはないが、非構造データリージョンはデータリージョンの開始と終了を異なるプログラムスコープで書くことができる。C++では、コンストラクタで配列Vを作り、ディストラクタでそれを消す。従って、図5に示すように、コンストラクタで配列Vのデータリージョンを開始し、ディストラクタでデータリージョンを終了するしか書きようが無い。しかし、これは一つのプログラム階層には収まらず、構造のないデータリージョンの定義が必要になるというわけである。

図5 C++では、コンストラクタでデータリージョンを開始(enter)し、ディストラクタで終了(exit)するように書くことが必要

また、多重の階層がある複雑なプログラムの場合は、どの{ }にデータリージョンが対応しているのかが分かりにくくなる。このような場合は、enter dataを書いておけば、その配列を含むデータリージョンが開始されるので、安心である。なお、enter dataは複数回行ってもよい。しかし、exit dateを行うとデバイス上のオブジェクトが消去されてしまうので、exitは1回しかできない。

これらのPragmaでのデバイスメモリの確保とデータのコピーに関する指定は、Data APIとして関数として呼び出すこともできるようになっている。

図6 これらのPragmaでの指定はAPIとして利用することもできる

OpenACCとして残っている大きな問題は、配列の要素がポインタを含み、ポインタが別のオブジェクトクトを指しているというケースが処理できないことである。この場合は、元の配列をデバイスメモリにコピーしても、ポインタが指しているオブジェクトがデバイスメモリにコピーされていなければ、GPUのカーネルプログラムでは目的のオブジェクトにアクセスすることはできない。

このように単純に配列をコピーするだけでは済まないケースをdeep copyと呼ぶ。OpenACC 2.0ではdeep copyは自動的には処理できないが、このData APIを使って、ポインタの先のオブジェクトをコピーして、ポインタを付け替えるという処理を記述することは可能である。