OpenACC(Open Acceleratorsの略)は、異種混合環境における並列コンピューティングを容易にするための標準規格です。主に
CPUとGPUが混在するシステムでの並列プログラミングを簡素化することを目的に、クレイ、CAPS、
NVIDIA、そしてPGIによって共同開発されました。
OpenACCは、
CUDAやOpenCLといったより低レベルな並列プログラミングAPIを抽象化したものであり、OpenMPと似た立ち位置にあります。プログラマーは、OpenMPと同様に、ディレクティブと追加のライブラリ関数を用いて、高速化したい箇所をC、
C++、またはFortranの
ソースコード上で指示します。OpenMP 4.0以降と同様に、OpenACCは
CPUとGPUの両方をターゲットにでき、これらのプロセッサ上で計算コードを実行できます。
OpenACCの標準化に関わるメンバーは、OpenMPの将来のバージョンでアクセラレータ(GPUや
コプロセッサ)をサポートすることや、OpenMPを拡張するための共通仕様を策定することを目指し、OpenMP標準化グループにも参加して活動しています。これらの活動は技術報告書にまとめられており、その議論は年次スーパーコンピューティング会議(2012年11月、
ソルトレイクシティ)で発表されました。この会議では、
NVIDIA以外のアクセラレータを製造するハードウェアベンダーの意見をOpenMPに反映させるという意思も示されました。
2012年の国際スーパーコンピューティング会議(ISC'12)では、OpenACCが
NVIDIA、AMD、Intelのアクセラレータ上で動作することが実証されました。ただし、具体的な性能データは公開されていません。2012年11月12日には、スーパーコンピューティング会議(SC12)でOpenACCバージョン2.0の草案が発表されました。このバージョンでは、データ転送の制御(非構造化データの取り扱いや不連続メモリへの対応の改善)、明示的な関数呼び出し、分割コンパイル(高速化されたコードライブラリの作成と再利用を可能にする)といった新機能が提案され、2013年1月に正式リリースされました。その後、バージョン2.5(2015年10月)、2.6(2017年11月)、2.7(2018年11月)と、継続的に仕様が更新されています。しかし、2019年4月3日には、クレイが同社のコンパイラCCE/9.0におけるOpenACCのサポートを終了することを発表しました。
コンパイラサポート
OpenACCは、PGI(コンパイラバージョン12.6以降)やクレイ(自社ハードウェアのみ)といった商用コンパイラでサポートされています。また、オープンソースのコンパイラも複数存在します。
OpenUH: CとFortranをサポートするOpen64をベースにしたコンパイラで、ヒューストン大学のHPCToolsグループが開発しています。
OpenARC: オークリッジ国立研究所が開発したCコンパイラで、OpenACC 1.0の仕様をほぼすべてサポートしています。
accUL: ラ・ラグーナ大学が開発している実験的なオープンソースコンパイラです。
Omni Compiler: 筑波大学と
理化学研究所が共同開発しているコンパイラで、OpenACCの他にXcalableMP、XcalableACCもサポートしています。
IPMACC: ビクトリア大学が開発しているCコンパイラで、OpenACCをCUDA、OpenCL、ISPCに変換します。現在のところ、data、kernels、loop、cacheといったディレクティブのみをサポートしています。
GCC(GNU Compiler Collection)へのOpenACCのサポートは、開発の進捗が遅れていました。2013年9月にサムスンがGPUをターゲットとしたOpenACC 1.1の実装を発表しましたが、これはOpenCLに変換するものでした。その後、NVIDIAがOpenACC 2.0をベースとした「真」の実装を発表しましたが、これはNVIDIA独自のPTXアセンブリ言語のみを対象としていたため、議論を引き起こしました。GCCバージョン5.1ではOpenACC/PTXの実験的サポートが終了しましたが、GCC 6とGCC 7ではOpenACC 2.0a仕様の改良版が搭載され、GCC 9.1ではOpenACC 2.5のサポートがほぼ完了したと発表されています。
使用方法
OpenACCでのプログラミングの基本は、OpenMP 3.xやOpenHMPPに似ており、ディレクティブを多用します。また、ランタイムライブラリには、いくつかのサポート関数が定義されています。C言語では`openacc.h`を、Fortranでは`openacc_lib.h`をインクルードし、`acc_init()`関数を呼び出す必要があります。
ディレクティブ
OpenACCは、幅広いディレクティブ(プラグマ)を定義しています。以下はその例です。
`#pragma acc parallel` および `#pragma acc kernels`
アクセラレータ上で実行される並列計算カーネルを定義します。
`#pragma acc data`
アクセラレータとCPU間のデータ転送を制御します。
`#pragma acc loop`
`parallel`または`kernels`領域内の並列性を定義します。
ランタイムAPI
OpenACCには、以下のようなランタイムAPI関数も定義されています。
`acc_get_num_devices()`, `acc_set_device_type()`, `acc_get_device_type()`, `acc_set_device_num()`, `acc_get_device_num()`
`acc_async_test()`, `acc_async_test_all()`, `acc_async_wait()`, `acc_async_wait_all()`
`acc_init()`, `acc_shutdown()`
`acc_on_device()`
`acc_malloc()`, `acc_free()`
OpenACCは、通常、ターゲットデバイス上で作業を組織化しますが、`gang`と`worker`を使って再定義することも可能です。`gang`は複数の`worker`から構成され、多数の処理要素を操作します(OpenCLのworkgroupに相当します)。
関連項目
XcalableACC
XcalableMP
C++ AMP
CUDA
OpenCL
OpenHMPP
OpenMP
出典
OpenACC 公式サイト
NVIDIAによるOpenACCの使用例
NVIDIAによるOpenACCの使用例
成瀬彰:「OpenACC で始めるGPUコンピューティング:OpenACC 概要」
成瀬彰:「OpenACCで始めるGPUコンピューティング:ループの並列化」
成瀬彰:「OpenACCで始めるGPUコンピューティング:データ転送の最適化」
成瀬彰:「OpenACCで始めるGPUコンピューティング:ループの最適化」
HPC WORLD / OpenACCとNVIDIA HPC SDKのユーザーコミュニティ
OpenACC ディレクティブによるプログラミング
「OpenACCとMPIによるマルチGPUプログラミング入門」(2020年12月22日、東京大学情報基盤センター講習会資料)
「OpenACCとMPIによるマルチGPUプログラミング入門」(2023年6月30日、東京大学情報基盤センター講習会資料)
Stéphane Ethier: Introduction to GPU programming with OpenACC
GPU移行
(RIST主催の)HPCプログラミングセミナーで使用する資料の公開ページ ※ OpenACCについては「アクセラレータ入門」の資料を参照。
「GPUプログラミング入門」(東京大学情報基盤センター,第188回お試しアカウント付き並列プログラミング講習会,2022年10月5日資料)