CUDAプログラムからOpenACCプログラムを呼び出す方法

CUDAプログラムからOpenACCプログラムを呼び出す

また、図19に示すように、#pragma acc parallel loopディレクティブにdeviceptr節を付けることができる。この例ではx、yはcudaMallocでGPUメモリ上に確保されている。そして、deviceptrは、xとyは既にデバイスメモリのアドレスなので、コンパイラとしては何も加工する必要はなく、GPUメモリ上のポインタとしてカーネル化したループに渡せば良いことを示している。

図20はCUDAで書いたメインから、OpenACCで書いたSaxpyを呼び出す例で、CUDA側では普通にx、yのメモリをデバイスメモリに確保し、saxpyを呼び出している。一方、OpenACCで書かれたsaxpyでは#pragma acc kernelsディレクティブにdeviceptr(x[0:n],y[0:n])を付け加えて、x、yがGPUメモリ上に確保されており、カーネルの中ではそのままポインタとして扱えばよいことを示している。なお、Cの場合はx、yの領域のサイズが分からないので、deviceptrのところで[0:n]を付けてサイズを明示している。

このように、host_data use_deviceを使えば、OpenACCのメインからCUDAの関数を呼び出すことができ、deviceptrを使えばCUDAメインからOpenACCの関数を呼び出すことができる。

GPUカーネルからGPUカーネルの呼び出しを可能にするルーチンディレクティブ

OpenACC 1.0では、GPUで実行されるカーネルから、他のGPUカーネルを呼び出すことができず、その場合は呼び出されるカーネルは、インライン展開して呼び出し元のカーネルに含めておくことが必要とされていた。それがOpenACC 2.0では、可能になった。しかし、最適化のためには、その関数がどのような並列度で実行されるのかを指定することが望ましい。

この指定には、#pragma acc routineディレクティブが使われる。この後に、どのような並列度で実行するのかを指定するgang、worker、vectorという指定が続く。

NVIDIAのCUDAは、スレッドの1～3次元の配列であるスレッドブロックとスレッドブロックの1～3次元の配列であるグリッドという単位でGPUに実行を依頼する。

OpenACCのgangはスレッドブロックのまとまりに対応する。#pragma acc kernels loop gang(100)と書くと、次の行から始まるカーネルを100個のスレッドブロックで実行するコードが作られる。そして、この後にvector(256)を付けると、256スレッドをまとめて一つのベクタとして同期して実行するというコードとなる。NVIDIAのGPUでは32スレッドをワープと呼び、これが最小の実行単位である。従って、ベクタのサイズは32の倍数とするのが無駄がない。

Workerはベクタの集まりで、gangはworkerの集まりであるが、worker指定は無視しているコンパイラもあり、通常は使わなくてもよいとLarkin氏は言う。

そして、seqは、このGPUカーネルは複数のスレッドから呼び出された場合にも、シーケンシャルに実行されることを指定する。全スレッドから共通に使用される変数を操作するようなカーネルはseqで実行するようにすれば安全である。

アトミックなメモリアセスを行うディレクティブ

OpenACCで作られるカーネルプログラムはマルチスレッドであるので、共通変数への複数のスレッドからのアクセスが入り乱れると結果がおかしくなってしまう。このため、共通変数の読み出しから更新して書き戻すまでの間は、別のスレッドはその変数にアクセスできないようにする必要がある。

#pragma acc atomicディレクティブを指定すると、次の｛ ｝の中の文は、他のスレッドのアクセスを排除して、一時には一つのスレッドだけがアクセスできるようになる。Atomicに続いて、read、write、update、captureが指定でき、updateが一般的なread-modify-writeに対応する。Caputureは、updateに加えてその変数をその後の処理で使えるようにする。ただし、OpenACC 2.0規格を見ると、captureの場合は、v=cnt++のような代入文となっており、図22の記述が正しいのかどかどうか確認できない。なお、何も付けないとデフォールトではupdateになる。

このようにOpenACC 2.0では、きめこまかな指定をして、高性能なGPUカーネルを生成するための機能が充実してきており、ユーザ側の関心も高まってきている。