OpenACC 2.0でサポートされたACC_MAP_DATA

ホストとデバイスメモリの対応を切り替えるmap_data

OpenACC 2.0ではACC_MAP_DATAという機能がサポートされた。図7では、cudaMallocでx_dという領域をデバイスメモリ上に確保し、それをacc_map_dataでOpenACCのxという配列にマップしている。y_dも同様にyという配列にマップしている。

そして#pragma acc parallel loopの中でx、yを使っているが、この部分がコンパイラでGPUカーネルに変換されるときには、xはx_dに、yはy_dに変換されて使用されることになる。

また、acc_unmap_dataという関数もあり、xとx_dの結合を切ることもできる。

一般に、GPUのメモリはCPUのメモリより小さいので、GPUでの処理には、部分的なデータを転送して、順に処理していくという方法がとられる。この時、CPU側の全データの中から、処理を分担する部分だけを抜き出し、同じデバイスメモリアドレスに対応させれば、同じ処理カーネルが使える。

普通のデータリージョンでは、ホストメモリの配列と、デバイスメモリの配列の対応は固定されているが、acc_map_data関数を使うと、この対応を変えることができ、ホストメモリの中から、GPUに処理させる部分を抜き出して、デバイスメモリの同じアドレスにマップすることができる。このようなやり方は、対象となるデータを覗く窓が次々のスライドしていくので、スライディングウインドウと呼ばれるが、図8に見られるように、acc_map_data関数を使うと記述が容易になる。そしてupdateを使って必要な部分だけのホストとデバイスメモの内容を一致させるようにすれば効率が良い。

この機能は、ダブルバッファを実装するときにも便利に使える。