色々な機能があるデータリージョン

GTC 2014において、NVIDIAのJeff Larkin氏は、OpenACC入門に引き続いて、「Advanced　Accelerated Computing using Compiler Directives」と題してOpenACCの上級編のチュートリアルを行った。

以下のレポートはLarkin氏の講義をもとに、筆者が再構成し説明などを加えたものである。もし、誤りがあるとすれば、それは筆者の理解不足によるものであり、Larkin氏の責任ではない。

データの移動を減らすデータリージョン

入門編で、データがGPU(デバイス)メモリに存在する期間をデータリージョンとして定義し、不要なデータ転送を省くことにより、処理性能を大きく改善できることを見てきた。このデータリージョンには、まだまだ、色々な機能がある。

図2は入門編で出てきた、最初にデータリージョンを定義し、copy(A)で、Aは開始時にCPUメモリからデバイスメモリにデータを転送し、終了時にCPUメモリに書き戻すことを指定し、Anewはデバイスメモリに領域を確保するがCPUメモリとの転送は行わないと指定する。

この処理では、これで良いのであるが、大きなプログラムになると、幾つものデータリージョンが必要になる。しかし、GPUにメモリ領域を確保する処理は、結構重い処理で時間がかかる。

このため、GPUメモリがひっ迫しているのでなければ、大きなデータリージョンを最初に作って、一度確保したメモリは使い続ける方が効率が良い。しかし、図2の書き方では、CPU(ホスト)メモリとのデータ転送はデータリージョンの開始時と終了時にしか行われないので、途中でデータのやり取りを必要とする場合にはうまく行かない。

これを可能にするのが、図3の#pragma acc update self( )と#pragma acc update device( )ディレクティブである。なお、図3はFortranの例なので、#pragma accではなく!$accと書かれている。

GPUで何かの処理を行ったら、内容を変更した配列をupdate selfでホストメモリに反映させ、CPUで何かの処理を行ったら、update deviceでデバイスメモリにコピーして内容を一致させられる。さらにupdateは配列Aの全体のデータをコピーするのではなく、配列Aの中のコピーする部分を指定することができる。配列Aの変更された部分が限定されている場合は、必要な部分だけをコピーするように指定すれば、転送するデータ量を減らすことができる。OpenACC 2.0でCPUメモリへの転送をupdate selfにしたのは、データリージョンが階層的に作れる仕様になり、デバイスからのデータがコピーされる先がupdateが書かれたデータリージョンの配列になったからである。