OpenACCが備える大きなメリットとは?

OpenACCのインタオペラビリティ

OpenACCは、ディレクティブの追加で比較的簡単にGPUプログラムを作ることができるだけでなく、CUDAやOpenCLと一緒に使うこともできる。すでにCUDAで書かれたプログラムがあり、新たに機能を付け加える場合、追加部分だけOpenACCで書くという使い方や、その逆にOpenACCのプログラムにCUDAで書いたコードを加えることもでき、最適化されたCUDAライブラリを使うこともできる。OpenACCを使う場合、全部をOpenACCで書き直す必要はない。これはOpenACCの大きなメリットである。

図15は、OpenACCのasyncキューとCUDAのstreamを繋ぐやり方を示したもので、acc_get_cuda_stream関数は、OpenACCのasyncキューの識別子を引数としてCUDAのストリームのハンドルに変換する。CUDAストリームのハンドルは、cublasなどに引数として渡して処理を行わせることができる。一方、acc_set_cuda_streams関数は、すでに作ってあるCUDAストリームを引数として渡したasyncキューに対応させる関数である。

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

図16のHost Dataディレクティブは、デバイスメモリのアドレスをホスト側のコードに記述できるようにする。#pragma acc data copy(x,y)でデバイスメモリにx、yのデータを入れるメモリ領域が確保されるが、#pragma acc host_data use_device(x,y)に続く｛ ｝の中ではCPUのx、yではなく、GPU側のx、yのアドレスが使われる。ある意味、ホスト側のx、yのアドレスをGPU側のアドレスに変換するように働く。

ホスト側のコードでデバイスメモリを直接アクセスするわけではないが、次の図16の例のように、host_data use_device( )を指定してx、yはデバイスのアドレスであることを指定する。そして、CUDAで書かれたsaxpyを呼び出している。右の下側に書かれたsaxpyでは、x、yはデバイスアドレスであることを示すためにdx、dyと書かれている。そしてsaxpyはsaxpy_kernelのグリッドを呼び出しているが、一貫してデバイスアドレスを使っているので、この間、アドレスの変換は必要なく、効率よく処理ができる。

図18では#pragma acc host_data use_device(x,y)に続いてcublasSaxpy(N,2.0,x,1,y,1)でCUDAに最適化したBLASライブラリであるcublasを呼び出している。