PGIトップ › PGI技術情報・TIPS › コンパイラ・オプションの使用方法 › PGI OpenACC 用のオプション

PGI OpenACC コンパイル用のオプション

対象 PGI アクセラレータ、CUDA Fortran

OpenACC

「アクセラレータ」とは、特別の目的で CPU にアタッチして使用する協調プロセッサであり、時間の掛かる計算部分を CPU の演算機構からデータと実行部分のカーネルをオフロードするために使用されます。ここでは、 GPUをアクセラレータとして使用する際のPGI アクセラレータ™ コンパイラのオプションについて説明します。また、PGI 12.6 からディレクティブベースの標準プログラミングモデルである OpenACC 1.0 に準拠し、現バージョンでは OpenACC 2.6 を包括的にサポートしております。OpenACC を使用するためのオプションも説明します。なお、OpenACC プログラミング・ガイドは、こちらにありますPGI 19.3 用更新 2019/3/28Copyright © 株式会社ソフテック

PGI OpenACC 用コマンド・オプション

 PGI の Fortran, C, C++ のコンパイラを使用して、PGI OpenACC 用ディレクティブを認識するためのオプションの例を示します。以下は、pgfortran(pgf95、あるいは、pgf90 も同じコンパイラです)を使用した場合の例ですが、C/C++ 言語用の pgcc/pgc++(pgcpp for Windows) コンパイラのオプションの設定方法も、同様です。なお、コマンドライン上でリンク時にも、必ず、コンパイル時に指定したものと同じ -acc 並びに-ta オプションを指定することが必要です。なお、この -ta=tesla オプションは、従来の PGI Accelerator Programming model 用のディレクティブ(OpenACC の前身)も認識します。但し、今後は、OpenACC を利用されることを推奨します。ここでは、OpenACC 用のコンパイル・オプションを説明します。なお、PGIコンパイラでは、両方のディレクティブ形式を認識しますが、注意すべきことは、一つのアクセラレータ領域の中で、PGI Accelerator 用ディレクティブと OpenACC 用のディレクティブの両方を使用することは許されません。どちらかの形式に統一して使用するようにして下さい。

 PGI 12.6以降、OpenACC 専用のコンパイル・リンク用オプションとして -acc が提供されました。これは、OpenACCディレクティブの認識を行います。さらに、これに加え、target accelerator の特性を指定するために、-ta オプションを指定することもできます。なお、PGI 14.1 以降、-ta オプションのサブオプションが一部変更となっておりますので、ご注意ください。

 PGI 15.1 以降、64-bitコンパイラのアクセラレータ用デバイスコードの生成プログラムが、デフォルトで nvvm を使用するように変更されました。今まで使用していた古い CUDA-C コード生成ソフトウェアを使用したい場合、あるいはデフォルトの llvm によるexecutable の実行に問題がある場合は、
-ta=tesla:nollvm あるいは -Mcuda=nollvm オプションを指定してください。

 PGI 17.1 以降Pascal(CC 6.0) GPU 用と過去の CC 2.0 用の executable はデフォルトで生成されません。これらの GPU 用の executable の生成には、-ta=tesla,cc60 or -ta=tesla,cc20 を付けてコンパイル、リンクしてください。

 PGI 17.7 以降Pascal(CC 6.0), Volta(CC 7.0) GPU 用と過去の CC 2.0 用の executable はデフォルトで生成されません。これらの GPU 用の executable の生成には、-ta=tesla,cc60,cuda9.0 or -ta=tesla,cc70,cuda9.0 or -ta=tesla,cc20 を付けてコンパイル、リンクしてください。Volta GPU をターゲットにする場合は、必ず CUDA 9 toolkit(-ta=tesla,cuda9.0)以上を使用する必要があります。

 PGI 18.7 以降CC 2.0 用の executable は生成されません。CC3.0~7.0用のコードを生成できます。

● Fortranコンパイルの一例
  pgfortran -fast -Minfo -acc test.f
    -acc の他に、-ta=tesla のサブオプションを付加することも可能
  pgfortran -fast -Minfo=accel -acc -ta=tesla,cuda10.0,cc70 test.f90 (cc70用 executable)

 ● C コンパイルの一例
  pgcc -fast -Minfo -acc test.c
    -acc の他に、-ta=tesla のサブオプションを付加することも可能
  pgcc -fast -Minfo=accel -acc -ta=tesla,keep test.c 

    なお、PGI 13.10 以前は、-ta=tesla の代わりに -ta=nvidia が有効です。
  • -Minfo=accel:このオプションを指定すると、コンパイラがアクセラレータ領域をGPUカーネルに翻訳できたかどうかについて、コンパイラのメッセージとして出力します。-Minfo のみの指定では、その他の最適化情報も併せてメッセージとして出力します。
  • -acc[=[no]autopar|[no]required|strict|verystrict|sync|[no]wait]:OpenACC ディレクティブを認識するオプションです。コンパイル時だけではなく、リンク時にも必要です。


    -acc=[no]autopar は、OpenACC parallel 構文内の自動並列化を行う[行わない]を指定します。(PGI 13.6以降)
    -acc=[no]required はアクセラレータ・コードを生成出来なかった場合、コンパイルエラーとする(default) (PGI 14.1 以降、但し、PGI 15.1 で廃止)
    -acc=strict は、non-OpenACC accelerator ディレクティブが見つかった場合、warning を出します。
    -acc=verystrict は、non-OpenACC accelerator ディレクティブが見つかった場合、エラーメッセージを出し、コンパイルを終了します。
    -acc=sync は、async cluase を無視します。
    -acc=[no]wait は各デバイス kernel の終了を待つか待たないかを指示します。

  • PGI 14.1 以降、-ta オプションの指定方法が変更されました。 デフォルトは -ta=tesla,host です。

    -ta=tesla(,tesla_suboptions) : NVIDIA 社の GPU ボード対応の OpenACC コンパイルを行います。従来の -ta="nvidia" の "nvidia"名 は廃止されました。

    -ta=multicore : ホスト上のマルチコアCPU上で並列動作するように OpenACC プログラムをコンパイルします。(PGI 15.10 以降) 実行に使用する CPU コア数は、ACC_NUM_CORES 環境変数に値をセットする。

    (PGI 13.10 以前の指定方法) -ta=nvidia(,nvidia_suboptions) : PGI アクセラレータディレクティブあるいは、OpenACC ディレクティブを認識するオプションです。-ta は、ターゲット・アーキテクチャを意味します。PGI 13.10 までは、「nvidia」 のみとなります。Fortran における !$acc ディレクティブ、C における #pragma acc ディレクティブをコンパイラに認識させ、ターゲットへの細かなオプションを使用するために、以下のサブオプションを使用できます。

  • -ta=tesla - NVIDIAアクセラレータをターゲットとして選択します。さらに、以下の tesla(nvidia) 用のサブオプションがあります。このサブオプションは、カンマ(,)で区切って複数のものを指定することができます。
サブオプション NVIDIA -ta=tesla(nvidia) のサブオプション
analysis ループの解析のみ行い、コードの生成を行いません。)(PGI 13.10以降廃止)
cc10 compute capability 1.0 のコードを生成 (PGI 14.1以降廃止)
cc11 compute capability 1.1 のコードを生成 (PGI 14.1以降廃止)
cc12 compute capability 1.2 のコードを生成 (PGI 14.1以降廃止)
cc13 compute capability 1.3 のコードを生成 (PGI 14.1以降廃止)
cc1x compute capability 1.x のコードを生成(PGI 15.1以降廃止)
cc1+ compute capability 1.x, 2.x, 3.x のコードを生成 (PGI 14.1以降)(PGI 15.1以降廃止)
tesla cc1x と同じ(PGI 13.1以降)(PGI 15.1以降廃止)
tesla+ cc1+ と同じ (PGI 14.1以降)(PGI 15.1以降廃止)
cc20 compute capability 2.0 のコードを生成 (PGI 10.4以降) (PGI 14.1以降廃止) (PGI 15.5以降復活)
cc2x compute capability 2.x のコードを生成 (PGI 10.4以降)
cc2+ compute capability 2.x, 3.x のコードを生成 (PGI 14.1以降)
fermi cc2xと同じ (PGI 13.1以降)
felmi+ cc2+と同じ (PGI 14.1以降)
cc30 compute capability 3.0 のコードを生成 (PGI 12.8以降) (PGI 14.1以降廃止) (PGI 15.5以降復活)
cc35 compute capability 3.5 のコードを生成 (PGI 13.1以降) (PGI 14.1以降廃止) (PGI 15.5以降復活)
cc3x compute capability 3.x のコードを生成 (PGI 12.8以降)
cc3+ compute capability 3.x (=cc3x) 以上のコードを生成 (PGI 14.1以降)
kepler cc3xと同じ (PGI 13.1以降)
kepler+ cc3+と同じ (PGI 14.1以降)
cc50 compute capability 5.0 (=cc50) (PGI 15.7以降)
cc60 compute capability 6.0 (=cc60) (PGI 16.9以降)
cc70 compute capability 7.0 (=cc70) (PGI 17.7以降)
cc75 compute capability 7.5 (=cc75) (PGI 18.10以降)
ccall 当該プラットフォームと、選択されたまたはデフォルトのCUDAツールキットによってサポートされているすべてのCCのためのコードを生成
charstring GPUカーネル内で文字列の使用を制限付きで使用する(PGI 15.1以降)
cuda2.3 or 2.3 PGIにバンドルされた CUDA toolkit 2.3 バージョンを使用 (PGI 10.4以降)
cuda3.0 or 3.0 PGIにバンドルされた CUDA toolkit 3.0 バージョンを使用 (PGI 10.4以降)
cuda3.1 or 3.1 PGIにバンドルされた CUDA toolkit 3.1 バージョンを使用 (PGI 10.8以降)
cuda3.2 or 3.2 PGIにバンドルされた CUDA toolkit 3.2 バージョンを使用 (PGI 11.0以降)
cuda4.0 or 4.0 PGIにバンドルされた CUDA toolkit 4.0 バージョンを使用 (PGI 11.6以降)
cuda4.1 or 4.1 PGIにバンドルされた CUDA toolkit 4.1 バージョンを使用 (PGI 12.2以降)
cuda4.2 or 4.2 PGIにバンドルされた CUDA toolkit 4.2 バージョンを使用 (PGI 12.6以降)
cuda5.0 or 5.0 PGIにバンドルされた CUDA toolkit 5.0 バージョンを使用 (PGI 13.1以降)
cuda5.5 or 5.5 PGIにバンドルされた CUDA toolkit 5.5 バージョンを使用 (PGI 13.9以降)
cuda6.0 or 6.0 PGIにバンドルされた CUDA toolkit 6.0 バージョンを使用 (PGI 14.4以降)
cuda6.5 or 6.5 PGIにバンドルされた CUDA toolkit 6.5 バージョンを使用 (PGI 14.9以降)
cuda7.0 or 7.0 PGIにバンドルされた CUDA toolkit 7.0 バージョンを使用 (PGI 15.4以降)
cuda7.5 or 7.5 PGIにバンドルされた CUDA toolkit 7.5 バージョンを使用 (PGI 15.9以降)
cuda8.0 or 8.0 PGIにバンドルされた CUDA toolkit 8.0 バージョンを使用 (PGI 16.10以降)
cuda9.0 or 9.0 PGIにバンドルされた CUDA toolkit 9.0 バージョンを使用 (PGI 17.9以降)
cuda9.1 or 9.1 PGIにバンドルされた CUDA toolkit 9.1 バージョンを使用 (PGI 18.1以降)
cuda9.2 or 9.2 PGIにバンドルされた CUDA toolkit 9.2 バージョンを使用 (PGI 18.5以降)
cuda10.0 or 10.0 PGIにバンドルされた CUDA toolkit 10.0 バージョンを使用 (PGI 18.10以降)
cuda10.1 or 10.1 PGIにバンドルされた CUDA toolkit 10.1 バージョンを使用 (PGI 19.3以降)
[no]debug デバイスコード内にデバッグ情報を生成する[しない] (PGI 14.1 以降)
deepcopy OpenACC Fortran における派生型変数の full deep copyを有効にする(PGI 17.7以降)
fastmath fast mathライブラリを使用
[no]flushz GPU上の浮動小数点演算の flush-to-zero モードを制御。デフォルトはnoflushz。 (PGI 11.5以降)
[no]fma fused-multiply-add命令を生成する[しない](-O3ではデフォルト)
keep kernelバイナリファイル(.bin)、kernelソースファイル(.gpu)、portable assembly(.ptx)ファイルを保持し、各々ファイルとして出力する。-ta=tesla,nollvmと共に使用する (PGI 13.10以降)
keepbin kernelバイナリファイルを保持し、ファイル(.bin)として出力する)(PGI 13.10以降廃止)
keepgpu kernelソースファイルを保持し、ファイル(.gpu)として出力する (PGI 13.10以降廃止)
keepptx GPUコードのためのportable assembly(.ptx)ファイルを保持し、ファイルとして出力する)(PGI 13.10以降廃止)
[no]lineinfo GPU line informationを生成する(PGI 15.1以降)
[no]llvm llvmベースのバックエンドを使用してコードを生成する。(PGI 15.1以降) 64-bit上ではLLVMバックエンドを使う [使わない]。なお、PGI 15.1 以降はデフォルト llvm を使うように変更された。リンク時にエラーが生じる場合、nollvm を試すことをお勧めする。
managed OpenACCの 割り当て可能なデータが CUDA Unified Memory に配置されている場合、明示的なデータ移動またはデータ指示を必要としない機能を有効にする。(PGI 17.7以降)
maxregcount:n GPU上で使用するレジスタの最大数を指定。ブランクの場合は、制約が無いと解釈する
mul24 添字計算に、24ビット乗算を使用 (GT200系、CC 1.3のみ)(PGI 13.10以降廃止)
noL1 グローバル変数をキャッシュするためのハードウェア L1 データキャッシュの使用を抑止する (PGI 13.10以降)
loadcache:L1
loadcache:L2
グローバル変数をキャッシュするためにハードウェア L1 あるいは L2データキャッシュを使用する。但し、アーキテクチャ上、有効とならないGPUがある (PGI 14.4以降)
pin デフォルトを pin ホストメモリ(割付)としてセットする(PGI 14.1~PGI 15.10)
pinned デフォルトを pin ホストメモリ(割付)としてセットする。プログラムのアロケート時に pinned メモリを割り付けるように変更した。(PGI 16.1以降)
time アクセラレータ領域の単純な時間情報を集積するためにプロファイル・ライブラリをリンクする。このオプションは、PGI 13.1 以降廃止されました。この代わりに、プロファイルを環境変数 PGI_ACC_TIME に 1 をセットすることにより実行後プロファイル情報が出力されます。
[no]required アクセラレータ・コードを生成出来なかった場合、コンパイルエラーとする[しない] (default)(PGI 15.1 以降廃止)
[no]rdc 異なるファイルに配置されたデバイスルーチンをそれぞれ分割コンパイルし、リンクが出来るようにする。cc2x以降、CUDA 5.0 以降の機能を使用する。(PGI 13.1以降 + CUDA 5.0 以降) (PGI 14.1 は以降デフォルト)
[no]unroll 自動的に最内側ループのアンローリングを行う (default at -O3)(PGI 14.9以降)
managed CUDA managed Memory を使用する
beta ベータ版機能のコード生成(生成コード内の 128-bit ロード・ストアオペレーションを有効化)(PGI 15.7以降)
[no]wait ホスト側での実行継続を行う際に、各カーネルが終了するまで待つ。nowaitは待たない。)(PGI 13.10以降廃止)
safecache cache directive 内での可変長の配列セクションの使用を許す。但し、そのサイズは CUDA shared memory 内に収まるものでなければならない。(PGI 16.5以降)
zeroinit 割付されるデバイスメモリの値をゼロに初期化する(PGI 18.1以降)
  • -ta=tesla(radeon),host - ターゲットとして、hostを選択する。アクセラレータ領域をホスト側で実行するようにコンパイルする。このオプションは、GPU が実装されていないシステムでも動作するような実行バイナリとなる PGI Unified Binaryコード を生成します。この host サブオプションは、上記の tesla 専用のサブオプションを指定した後に指定します(一番最後に指定します)
  • (更新情報:PGI 10.3以降) GPUデバイスの Compute capability の指定オプション -Mcuda=cc?? を複数指定することができます。PGI 10.2 以前のデフォルトのGPUデバイス Copmute capability は、1.3(cc13) でした。PGI 10.3 以降、デフォルトのターゲットは、1.0(cc10) 並びに 1.3(cc13) となりました。さらに、複数の GPU compute capability のターゲットとするようなコード生成したい場合は、-Mcuda=cc10 -Mcuda=cc11 -Mcuda=cc12 -Mcuda=cc13 という風に、コマンドラインで複数指定することができます。
  • (更新情報:PGI 10.4以降)CUDA 3.0 toolkitを使用してコンパイルした場合(-ta=nvidia:cuda3.0)、あるいは、コンパイラの初期化ファイル sitercファイルの中にset CUDAVERSION=3.0 を設定した場合は、2.3 CUDA ドライバ実装システム上では動作しませんのでご注意下さい。CUDAドライバも CUDA 3.0 用に変更する必要があります。現在、システムに実装されている CUDA ドライバを確認する方法は、pgaccelinfo を実行して下さい。
  • (更新情報:PGI 10.8以降)PGI 10.8 にて、CUDA 3.1 toolkit がバンドル・実装されました。これに伴い、CUDA toolkit 3.0 はバンドルされていません。なお、デフォルトでは CUDA tool kit (CUDA のライブラリ等)2.3 を使用しますので、3.1 toolkitを使用したい場合は、明示的に cuda3.1 のサブオプションを指定してコンパイル・リンクする必要があります。
  • (更新情報:PGI 11.0以降)PGI 11.0 にて、CUDA 3.2 toolkit がバンドル・実装されました。これに伴い、CUDA toolkit 3.1, 3.2 がバンドルされております。なお、PGI 11.0~11.5では、デフォルトでCUDA tool kit (CUDA のライブラリ等)3.1 を使用しますので、3.2 toolkitを使用したい場合は、明示的に cuda3.2 のサブオプションを指定してコンパイル・リンクする必要があります。
  • (更新情報:PGI 11.6以降)PGI 11.6 にて、CUDA 4.0 toolkit がバンドル・実装されました。これに伴い、CUDA toolkit 3.2, 4.0 がバンドルされております。なお、PGI 11.6 以降では、デフォルトでCUDA tool kit (CUDA のライブラリ等)3.2 を使用しますので、4.0 toolkitを使用したい場合は、明示的に cuda4.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 4.0ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 12.2以降)PGI 12.2 にて、CUDA 4.1 toolkit がバンドル・実装されました。これに伴い、CUDA toolkit 4.0, 4.1 がバンドルされております。なお、PGI 12.2 以降では、デフォルトでCUDA tool kit (CUDA のライブラリ等)4.0 を使用しますので、4.1 toolkitを使用したい場合は、明示的に cuda4.1 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 4.1ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 12.6以降)PGI 12.6 にて、CUDA 4.2 toolkit がバンドル・実装されました。これに伴い、CUDA toolkit 4.1, 4.2 がバンドルされております。なお、PGI 12.6 以降では、デフォルトでCUDA tool kit (CUDA のライブラリ等)4.1 を使用しますので、4.2 toolkitを使用したい場合は、明示的に cuda4.2 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 4.2ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 13.1以降)PGI 13.1 にて、CUDA 5.0 toolkit がバンドル・実装されました。これに伴い、CUDA toolkit 4.2, 5.0 の二つがバンドルされております。なお、PGI 13.1 以降では、デフォルトでCUDA tool kit (CUDA のライブラリ等)4.2 を使用しますので、5.0 toolkitを使用したい場合は、明示的に cuda5.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 5.0ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 13.9以降)PGI 13.9 にて、CUDA 5.5 toolkit が新たにバンドル・実装されました。これに伴い、CUDA toolkit 5.0, 5.5 の二つがバンドルされております。なお、PGI 13.9 以降では、デフォルトでCUDA tool kit (CUDA のライブラリ等)5.0 を使用しますので、5.5 toolkitを使用したい場合は、明示的に cuda5.5 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 5.5ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 14.4以降)PGI 14.4 にて、CUDA 6.0 toolkit が新たにバンドル・実装されました。これに伴い、CUDA toolkit 5.5, 6.0 の二つがバンドルされております。なお、PGI 14.4 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)5.5 を使用しますので、6.0 toolkitを使用したい場合は、明示的に cuda6.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 6.0ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 14.9以降)PGI 14.9 にて、CUDA 6.5 toolkit が新たにバンドル・実装されました。これに伴い、CUDA toolkit 6.0, 6.5 の二つがバンドルされております。なお、PGI 14.9 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)6.0 を使用しますので、6.5 toolkitを使用したい場合は、明示的に cuda6.5 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 6.5ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 15.4以降)PGI 15.4 にて、CUDA 7.0 toolkit が新たにバンドル・実装されました。これに伴い、CUDA toolkit 6.5, 7.0 の二つがバンドルされております。なお、PGI 15.4 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)6.5 を使用しますので、7.0 toolkitを使用したい場合は、明示的に cuda7.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 7.0ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 15.9以降)PGI 15.9 にて、CUDA 7.5 toolkit が新たにバンドル・実装されました。これに伴い、CUDA toolkit 7.0, 7.5 の二つがバンドルされております。なお、PGI 15.9 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)7.0 を使用しますので、7.5 toolkitを使用したい場合は、明示的に cuda7.5 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 7.5ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 16.10以降)PGI 16.10 にて、CUDA 8.0 toolkit が新たにバンドル・実装されました。これに伴い、CUDA toolkit 7.0, 7.5, 8.0 の三つがバンドルされております。なお、PGI 16.10 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)7.0 を使用しますので、7.5 or 8.0 toolkitを使用したい場合は、明示的に cuda7.5 or cuda8.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 8.0 ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 17.1以降)PGI 17.1 では、CUDA toolkit 7.5, 8.0 の二つがバンドルされております。なお、PGI 17.1 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)7.5 を使用しますので、8.0 toolkitを使用したい場合は、明示的に cuda8.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 8.0 ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 17.9以降)PGI 17.9 では、CUDA toolkit 7.5, 8.0, 9.0(RC) の三つがバンドルされております。なお、PGI 17.9 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)7.5 を使用しますので、8.0 or 9.0 toolkitを使用したい場合は、明示的に cuda8.0 or cuda9.0 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA 8.0 or 9.0 ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 18.1以降)PGI 18.1 では、CUDA toolkit 8.0, 9.0, 9.1 の三つがバンドルされております。なお、PGI 18.1 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)8.0 を使用しますので、9.0 or 9.1 toolkitを使用したい場合は、明示的に cuda9.0 or cuda9.1 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA9.0 or 9.1 ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 18.5以降)PGI 18.5 では、CUDA toolkit 8.0, 9.0, 9.1, 9.2 の四つがバンドルされております。なお、PGI 18.5 以降では、デフォルトで CUDA tool kit (CUDA のライブラリ等)8.0 を使用しますので、9.0 or 9.1 or 9.2 toolkitを使用したい場合は、明示的に cuda9.0 or cuda9.1 or cuda9.2 のサブオプションを指定してコンパイル・リンクする必要があります(但し、CUDA9.0 or 9.1 or 9.2 ドライバーを実装してあることが必要です)。
  • (更新情報:PGI 18.7以降)コンパイラは、コンパイルに使用するシステムにインストールされている CUDA ドライバと一致するように、デフォルトの CUDA バージョンを設定するようになりました。 CUDA 9.1 および CUDA 9.2 toolkit は、PGI 18.7 リリースにバンドルされています。 CUDA 8.0 や CUDA 9.0 を含む旧式の CUDA ツールチェーンは、サイズを最小限に抑えるために PGI インストールパッケージから削除されましたが、新しい共存機能を使用して引き続きサポートされています。 -ta=tesla,cudaX.Y サブオプションを指定しない場合の扱いは、使用するCUDA Toolkitのバージョン制御をご覧ください。
  • (更新情報:PGI 19.3以降)コンパイラは、コンパイルに使用するシステムにインストールされている CUDA ドライバと一致するように、デフォルトの CUDA バージョンを設定します。 CUDA 9.2 および CUDA 10.0, 10.1 toolkit は、PGI 19.3 リリースにバンドルされています。 CUDA 9.0 等の旧式の CUDA ツールチェーンは、サイズを最小限に抑えるために PGI インストールパッケージから削除されましたが、新しい共存機能を使用して引き続きサポートされています。 -ta=tesla,cudaX.Y サブオプションを指定しない場合の扱いは、使用するCUDA Toolkitのバージョン制御をご覧ください。
    $ pgaccelinfo
    CUDA Driver Version            10010
    (以下省略)
    ここで、
    2.3 driverの場合: CUDA Driver Version 2030
    3.0 driverの場合: CUDA Driver Version 3000 (PGI 10.4以降)
    3.1 driverの場合: CUDA Driver Version 3010 (PGI 10.8以降)
    3.2 driverの場合: CUDA Driver Version 3020 (PGI 11.0以降)
    4.0 driverの場合: CUDA Driver Version 4000 (PGI 11.6以降)
    4.1 driverの場合: CUDA Driver Version 4010 (PGI 12.2以降)
    4.2 driverの場合: CUDA Driver Version 4020 (PGI 12.6以降)
    5.0 driverの場合: CUDA Driver Version 5000 (PGI 13.1以降)
    5.5 driverの場合: CUDA Driver Version 5050 (PGI 13.9以降
    6.0 driverの場合: CUDA Driver Version 6000 (PGI 14.4以降
    6.5 driverの場合: CUDA Driver Version 6050 (PGI 14.9以降)
    7.0 driverの場合: CUDA Driver Version 7000 (PGI 15.4以降)
    7.5 driverの場合: CUDA Driver Version 7050 (PGI 15.9以降)
    8.0 driverの場合: CUDA Driver Version 8000 (PGI 16.10以降)
    9.0 driverの場合: CUDA Driver Version 9000 (PGI 17.9以降)
    9.1 driverの場合: CUDA Driver Version 9010 (PGI 18.1以降)
    9.2 driverの場合: CUDA Driver Version 9020 (PGI 18.5以降)
    10.0 driverの場合: CUDA Driver Version 1000 (PGI 18.10以降)と表示されます
    10.1 driverの場合: CUDA Driver Version 10010 (PGI 19.3以降)と表示されます
  • (CUDA Toolkit バージョンのデフォルト値の変更) PGIコンパイラの各リビジョンは、デフォルトで使用する CUDA toolkitのバージョンが予め指定されています。CUDA Fortran あるいは、PGI Acceleratorアプリケーションは、デフォルトでは、この CUDA Toolkit バージョン(上記の各リビジョンのToolkitの説明を参照してください)を利用してビルドします。なお、この使用するCUDA Toolkitのバージョンのデフォルト値の変更が可能です。コンパイラソフトウェアを実装している $PGI/{target}/{version番号}/bin 配下に siterc ファイルがありますので、その中に、以下を挿入します。ここのパスで言う {target} は、linux86-64, win64, linuxpower 等を指します。
    (一例)
    set DEFCUDAVERSION=9.2;
    
    上記の設定を行わず、コンパイラコマンド上のオプションで明示的に 9.2 Toolkit (PGI 18.5 以上の場合)を使用するように指示するには、以下のように指定します。
    pgfortran -ta=tesla,cuda9.2,.... 
    pgcc -ta=tesla,cuda9.2,....
    
  • PGI の各リビジョンで追加・変更されたオプションについては、最新の「リリースノート」をご覧下さい

 コンパイラは、自動的に必要とする CUDA ソフトウェアのツールを発動し、GPU カーネルコードを生成し、オブジェクト・ファイルの中にカーネルを埋め込みます。

使用するCUDA Toolkitのバージョン制御 (PGI 18.7以降)

 最新バージョンの「CUDA Toolkitのバージョン制御」に関しては、リリースノートに記述しています。以下の記述は、PGI 2018 を例にしたものです。

 PGIコンパイラは、NVIDIA GPUで実行するプログラムを構築する際に、NVIDIAのCUDAツールキットを使用します。 すべてのPGIインストールパッケージは、必要なCUDA Toolkitコンポーネントを2018/cudaというPGIインストールディレクトリに置かれます。

 NVIDIA CUDAドライバは、そのシステムでGPU用にコンパイルされたプログラムを実行する前に、GPUを搭載したシステムにインストールする必要があります。 PGI製品にはCUDAドライバが含まれていません。 NVIDIAから適切なCUDAドライバをダウンロードしてインストールする必要があります。 CUDAドライバのバージョンは、少なくともコードをコンパイルしたCUDA Toolkitのバージョンと同じかあるいは新しいものでなければなりません。

 PGIツール pgaccelinfo は、最初の出力行としてドライバのバージョンを表示します。 ご使用のシステムにどのバージョンのCUDAドライバがインストールされているかわからない場合に使用してください。

PGI 18.7 には、次のバージョンのCUDAツールキットが含まれています。

  • CUDA 9.1
  • CUDA 9.2

 コンパイラに、使用するバージョンのCUDAツールキットを選択させるか、特定のバージョンを使用するように指示することができます。 このセクションの残りの部分では、すべてのオプションについて説明します。
 CUDAツールキットのバージョンを指定しない場合、コンパイラは、コンパイルするシステムにインストールされている CUDA ドライバのバージョンを使用して、使用するCUDAツールキットを決定します。 この自動検出機能は、PGI 18.7 の新機能です。 自動検出は、同じシステムでアプリケーションをコンパイルして実行する場合に特に便利です。 これはどのように動作するのですか? 他の情報がない場合、コンパイラは、システムにインストールされている CUDA ドライバのバージョンと一致する PGI 2018/cuda ディレクトリ内の CUDA Toolkit バージョンを探します。 一致するものが見つからない場合、コンパイラは CUDA ドライバのバージョンより新しいものではない最新のCUDAツールキットのバージョンを検索します。 CUDAドライバがインストールされていない場合、PGI 18.7 コンパイラはデフォルトの CUDA 9.1 に戻ります。 いくつかの例を見てみましょう。
 インストールした唯一のPGIコンパイラがPGI 18.7ならば、

  • CUDA Driverが9.2の場合、コンパイラはCUDA Toolkit 9.2を使用します。
  • CUDA Driverが9.1の場合、コンパイラはCUDA Toolkit 9.1を使用します。
  • CUDA Driverが9.0の場合、コンパイラはCUDA Toolkit 9.0が見つからなかったというエラーを出します。 CUDA Toolkit 9.0はPGI 18.7にバンドルされていません。
  • コンパイルシステムにCUDAドライバがインストールされていない場合、コンパイラはCUDA Toolkitバージョン9.1を使用します。
  • CUDAドライバがCUDA 9.2より新しい場合、コンパイラは引き続きCUDAツールキット9.2を使用します。 コンパイラは、CUDAドライバより新しいものではないと判明した最新のCUDAツールキットを選択します。

 次のいずれかの方法を使用して、CUDA Toolkitバージョンのコンパイラのデフォルト選択を変更することができます。

  • コンパイラオプションを使用します。 cudaX.Y サブオプションを -Mcuda または-ta=tesla に設定します。XY は CUDA バージョンを表します。 たとえば、Cファイルを CUDA 9.2 Toolkitでコンパイルするには、次のようにします。
    pgcc -ta=tesla:cuda9.2
    コンパイラオプションを使用すると、コンパイラを1回呼び出すたびにCUDAツールキットのバージョンが変更されます。
  • rfile 変数を使用します。 DEFCUDAVERSION を定義する行を、インストールした際の PGI の bin/ディレクトリの siterc ファイル、またはホームディレクトリの .mypgirc という名前のファイルに追加します。 たとえば、CUDA 9.2 Toolkitをデフォルトとして指定するには、次のいずれかの行に次の行を追加します。
    set DEFCUDAVERSION=9.2;
    rcfile変数を使用すると、rcfile を読み取るコンパイラのすべての呼び出しに対して CUDA Toolkit のバージョンが変更されます。

 CUDA Toolkitバージョンを指定すると、現在のPGIコンパイラにバンドルされているデフォルトとは異なるCUDA Toolkitインストールを使用するようにコンパイラに指示することもできます。 ほとんどのユーザーは、PGIで提供されているもの以外のCUDA Toolkitインストールを使用する必要はありませんが、この機能が必要な状況が発生します。 プレリリースCUDAソフトウェアを使用している開発者は、PGIリリースに含まれていないCUDAツールキットのバージョンでテストする必要があることがあります。 逆に、PGIリリースでインストールされた最も古いCUDAツールキットよりも古いCUDAツールキットでコンパイルする必要がある開発者もいます。 これらのユーザのために、PGIコンパイラは、PGIインストールディレクトリの外にインストールされたCUDAツールキットのコンポーネントと相互運用できます。

 PGIは、一緒にインストールされたバージョンのCUDAツールキットを使用して広範にテストし、その使用を完全にサポートします。 PGIインストールに含まれていないCUDA Toolkitコンポーネントの使用は、機能の違いが存在する可能性があることを理解して行われます。
PGIコンパイラでインストールされたバージョン以外のCUDAツールキットでコンパイルする機能は、すべてのプラットフォームでサポートされています。 Windowsプラットフォームでは、この機能はCUDA Toolkitバージョン9.2以降でサポートされています。

 PGI 18.7の CUDA 8.0 など、PGIリリースとともにインストールされていないCUDAツールキットを使用するには、3つのオプションがあります。

  • 基本デフォルトを上書きするには、rcfile 変数 DEFAULT_CUDA_HOME を使用します。
    set DEFAULT_CUDA_HOME=/opt/cuda-8.0;
  • 環境変数 CUDA_HOME を設定する
    export CUDA_HOME=/opt/cuda-8.0
  • コンパイラのコンパイル行の割り当てを使用するCUDA_HOME=
    pgfortran CUDA_HOME=/opt/cuda-7.5
  • コンパイラのコンパイル行の割り当てを使用するCUDA_HOME =
    pgfortran CUDA_HOME=/opt/cuda-8.0

 PGIコンパイラは、使用する CUDAツールキットのバージョンを決定する際に、次の優先順位を使用します。

  1. コンパイラに使用するCUDAツールキットのバージョンを知らせない場合、コンパイラは、システムにインストールされているCUDAドライバのバージョンと一致するPGIインストールディレクトリ2018 / cudaからCUDAツールキットを選択します。 PGIインストールディレクトリに直接一致が含まれていない場合は、CUDAドライババージョンより新しいバージョンではないディレクトリ内の最新バージョンが使用されます。 システムにCUDAドライバがインストールされていない場合、コンパイラは内部デフォルトに戻ります。 PGI 18.7では、このデフォルトはCUDA 9.1です。
  2. rcfile 変数 DEFAULT_CUDA_HOME は、デフォルトのデフォルトを上書きします。
  3. 環境変数 CUDA_HOMEは上記のすべてのデフォルトを上書きします。
  4. 環境変数PGI_CUDA_HOMEは上記のすべてを上書きします。 すでに定義されているCUDA_HOMEをオーバーライドする必要がある場合に備えて、上級ユーザーが利用できます。
  5. -Mcudaと-ta=tesla のユーザ指定の cudaX.Y サブオプションは、上記のデフォルトをすべて上書きし、PGIインストールディレクトリ2018/cudaにあるCUDAツールキットが使用されます。
  6. コンパイラのコンパイル行の割り当て CUDA_HOME= は、上記のすべてのデフォルト(cudaX.Yサブオプションを含む)を上書きします。

デフォルトのCompute Capabilityの設定方法

 PGI 18.7 以降、コンパイラは、NVIDIA GPU compute capability 3.0〜7.0 用のコードを生成できます。 コンパイラは、コンパイル時のシステムで検出された GPU のコンピューティング機能と一致する compute capability をデフォルトリストとして作成します。GPU が検出されない場合、コンパイラは cc35、 cc50、 cc60、および cc70を選択します。
 コマンドラインオプションまたは rcfile を使用して、1つまたは複数のコンピューティング機能を指定することによって、デフォルトを上書きすることができます。
 コマンドラインオプションを使用してデフォルトを変更するには、OpenACC の場合は -ta=tesla: に、CUDA Fortran の場合は -Mcuda= にcompute capabilityのカンマ区切りリストを入力します。rcfileを使用してデフォルトを変更するには、インストールの bin ディレクトリにある siterc ファイルの DEFCOMPUTECAP 値を空白で区切ったcompute capability のリストに設定します。

set DEFCOMPUTECAP=60 70;

 また、siterc ファイルを変更する権限がない場合は、ホームディレクトリの別の.mypgircファイル(Windows では mypgi_rc)に DEFCOMPUTECAP 定義を追加できます。デバイスコードの生成には時間がかかることがあるため、コンピューティング機能の数が増えるとコンパイル時間が長くなることがあります。

PGIアクセラレータ用実行時の環境変数

  • PGI_ACC_DEVICE_TYPE :  PGI_ACC_DEVICE_TYPE 環境変数は、プログラムの実行モジュールが一つ以上の異なるデバイスタイプを使用して実行できるように生成されていた場合(PGI Unified Binary)、「アクセラレータ・リージョン」を実行する際に使用するデフォルトのデバイスタイプを指定するものです。この環境変数の値は、コンパイラ・リリースのおける実装時に定義されていますが、現在、NVIDIA(nvidia) と HOST(host) が定義されています。
    例:

    export PGI_ACC_DEVICE_TYPE=NVIDIA
    setenv PGI_ACC_DEVICE_TYPE NVIDIA

  • ACC_DEVICE_NUM : ACC_DEVICE_NUM 環境変数は、「アクセラレータ・リージョン」を実行する際に使用するデフォルトのデバイス番号を指定するものです。環境変数の値は、0~正の整数でなければなりません。システム内に複数のGPUデバイスが実装されている場合、その論理番号が 0 から順番に付されて管理されています。pgaccelinfo コマンドを実行すると、各GPUデバイスのロプロパティが論理番号順に表示できます。0 を指定した場合、システム実装時のデフォルトが使用されます。
    例:

    export ACC_DEVICE_NUM=1
    setenv ACC_DEVICE_NUM 1

  • PGI_ACC_NOTIFY - (PGI 14.1 以降新設)PGI_ACC_NOTIFY 環境変数は、ビットマスクとして利用する整数定数をセットして、デバイス上の実行イベントの情報を出力するためのものです。整数値 1 をセットすると Kernel launch のイベントを標準出力として出力します。整数値 2 は、データ転送のイベントの出力、整数値 4 の場合は、region の entry/exit 情報、整数値 8 は、デバイス上の wait/sync のイベントを出力します。0を指定するとこの機能を抑止します(デフォルト)。
    例:

    export PGI_ACC_NOTIFY=2
    setenv PGI_ACC_NOTIFY 2
    (メッセージ一例)
    upload CUDA data file=acc_f2a.f90 function=main line=37 device=0 variable=a bytes=4000000
    download CUDA data file=acc_f2a.f90 function=main line=41 device=0 variable=r bytes=4000000

  • PGI_ACC_TIME - PGI_ACC_TIME 環境変数は、実行後に簡易プロファイル情報を標準出力に出力するために使用されます。PGI 12.x まで有効であった -ta=time の "time" sub-option は廃止されました。環境変数の値は、負の整数であってはなりません。0を指定するとこの機能を抑止します(デフォルト)。0以外の正数の場合は、プロファイル情報を標準出力に印字します。(PGI 13.1 以降)なお、この機能を有効にするには、LD_LIBRARY_PATH環境変数に、PGIのライブラリパスを設定する必要があります。具体的には、64ビット環境では、$PGI/linux86-64/{バージョン番号}/libをセット、32ビット環境では、$PGI/linux86/{バージョン番号}/libをセットして、当該プログラムを実行します。
    例:

    export PGI_ACC_TIME=1
    setenv PGI_ACC_TIME 1

  • PGI_ACC_DEBUG - PGI_ACC_DEBUG 環境変数は、PGIランタイムに、デバイスのメモリ割り当て、データ移動、カーネルの起動などの情報を生成するように指示するために使用する。1 をセットする。 PGI_ACC_DEBUG は、ほとんどがランタイム自体のデバッグに使用するように設計されていますが、プログラムがデバイスとどのようにやりとりするかを理解するのに役立つ。
  • PGI_ACC_BUFFERSIZE - (PGI 14.1 以降新設)PGI_ACC_BUFFERSIZE 環境変数は、NVIDIA デバイスにおけるホストとデバイス間のデータ転送で使用される pinned buffer(Pinned memory上)のサイズを指定するものである。
  • PGI_ACC_CUDA_GANGLIMIT - (PGI 14.1 以降新設)PGI_ACC_CUDA_GANGLIMIT 環境変数は、NVIDIA デバイスにおける、kernel によって起動される gang(CUDA thread block) の最大数を指定するものである。
  • PGI_ACC_DEV_MEMORY - (PGI 14.1 以降新設)PGI_ACC_DEV_MEMORY 環境変数は、AMD Radeon デバイスにおける、アロケートされる OpenCL のバッファの最大値を指定するものである。この最大値は、ターゲット・デバイスによって制限される場合がある。
  • ACC_NUM_CORES - ACC_NUM_CORES 環境変数は、-ta=multicore オプションで作成された マルチコアCPU 上で動作する OpenACC プログラムを動作させる際の利用するコア数を指定するものです。
  • ACC_BIND - ACC_BIND 環境変数は、-ta=multicore オプションで作成された マルチコアCPU 上で動作する OpenACC プログラムを動作させる際、デフォルトでセットされます。 ACC_BIND の設定挙動は、OpenMP における MP_BIND の挙動と同様です。
  • PGI_ACC_SYNCHRONOUS - PGI_ACC_SYNCHRONOUS 環境変数は、非同期の launch やデータ転送処理を抑止するように指示するものです。

PGI 13.2 以降の Windows 版で、 PGI_ACC_TIME やACC_NOTIFY 環境変数をセットして機能を使用する場合は、こちらの FAQ で説明する方法で DLL ライブラリの検索パスをセットしてください。

OpenACC における CUDA Unified Memory 機能の使用

PGI 17.7リリースでは、allocatable として割り付けられたデータ用の CUDA Unified Memory の使用がベータ版からプロダクション版に移行しました。 この機能は、 OpenACC および CUDA Unified Memory PGInsider の記事で詳しく説明されており、Linux/x86-64 および Linux/OpenPOWER コンパイラで利用できます。 Linux/x86-64 では、デフォルトの PGI コードジェネレータと Beta版 LLVMベースのコードジェネレータの両方を使用してサポートされています。 この機能を有効にするには、オプション -ta=tesla:managed をコンパイラおよびリンカのコマンドラインに追加します。

-ta=tesla:managed の場合、プログラム単位内のすべての C/C++、Fortran の明示的な allocate 文は、CUDA Unified Memory 上にデータを配置する同等の「管理された」データ割り当て呼び出しに置き換えられます。 管理対象データは CPU/GPU の単一アドレスを共有し、CPU と GPU メモリ間のデータ移動は CUDAドライバによって暗黙的に処理されます。 したがって、OpenACC のデータ句とディレクティブは、「管理された」データには必要ありません。 それらは本質的に無視され、実際には省略することができます。

プログラムが管理されたメモリを割り当てると、デバイスのメモリだけでなくホストの固定メモリも割り当てられます。したがって、割り当てと解放の操作がやや高価になり、データの転送がやや高速になります。 メモリプールアロケータは、アロケートおよびフリーオペレーションのオーバーヘッドを軽減するために使用されます。プールアロケータは -ta= tesla:managed または -ta= tesla:pinned に対してデフォルトで有効になっています。 PGI 17.7リリースでは、 -Mcuda が存在するとプールアロケータが無効になります。 今後のリリースでその制限を解除する作業を進めています。

管理データのデータ移動は、NVIDIA CUDA GPUドライバによって制御されます。 CPU または GPU 上でデータにアクセスするたびに、最後にアクセスされたときに同じデバイス上にデータが転送される可能性があります。 場合によっては、ページのスラッシングが発生し、パフォーマンスに影響を与えることがあります。 Parallel Forall では、CUDA Unified Memory の概要を紹介しています。

この機能には次の制限があります。

  • 管理対象メモリの使用は、動的に割り当てられたデータにのみ適用されます。静的データ(C静的および外部変数、Fortranモジュール、共通ブロックおよび保存変数)および関数ローカルデータは、OpenACC ランタイムによって処理されるため、従来通り、ユーザによるデータ移動管理が必要です。 動的に割り当てられた Fortran ローカル変数と Fortran 割り当て可能配列は暗黙的に管理されますが、Fortran 配列ポインタは管理されません。
  • ローカル、グローバル、または静的データを指すメンバを持つ割り当て可能な集合体を指定すると、 -ta=tesla:managedでコンパイルし、計算カーネルからそのポインタを介してメモリにアクセスしようとすると、実行時にエラーが発生します。
  • C ++仮想関数はサポートされていません。
  • ファイルにOpenACCコードがない場合でも、変数が割り当てられたファイルをコンパイルするには、 -ta=tesla:managedコンパイラ・オプションを使用する必要があります。

この機能には、NVIDIA Kepler GPUで 使用するときに次の追加制限があります。

  • Kepler GPU 上のデータの動きは、高速の固定非同期データ転送によって実現されます。 しかし、プログラムの観点からは、転送は同期的です。
  • PGI runtime は、Keplar GPU を搭載したシステムで-ta= tesla:managed が使用されている場合、カーネルの同期実行を強制します。 このような状況は、余分な同期と CPU とGPU のオーバーラップの減少により、パフォーマンスが低下する可能性があります。
  • 管理されるメモリの総量は、Kepler GPU で利用可能なデバイスメモリの量に制限されています。

この機能は、NVIDIA Fermi GPU ではサポートされていません。

CUDA Unified Memory Pool Allocator

動的メモリ割り当ては、cudaMallocManaged() を使用して行われます。このルーチンは、cudaMalloc()を使用して非ユニファイドメモリを割り当てるよりもオーバーヘッドが高いルーチンです。 cudaMallocManaged() の呼び出し回数が増えるほど、パフォーマンスへの影響が大きくなります。

cudaMallocManaged() 呼び出しのオーバーヘッドを軽減するために、-ta=tesla:managed および -ta=tesla:pinned は CUDAUnified Memory プールアロケータを使用して、cudaMallocManaged() への呼び出し回数を最小限に抑えます。 プールアロケータは、デフォルトで有効になっています。 次の環境変数を使用して無効にするか、動作を変更できます。

プールアロケータ環境変数

PGI_ACC_POOL_ALLOC プールアロケータを無効にします。 プールアロケータはデフォルトで有効になっています。 無効にするには、PGI_ACC_POOL_ALLOCを0に設定します。
PGI_ACC_POOL_SIZE プールのサイズを設定します。 デフォルトのサイズは1GBですが、他のサイズ(2GB、100MB、500KBなど)を使用できます。 実際のプール・サイズは、フィボナッチ・シリーズのサイズが、提供されたサイズまたはデフォルトサイズと比較して、最も近い、より小さい数になるように設定されます。 必要であれば、プールアロケータはプールを追加しますが、PGI_ACC_POOL_THRESHOLD値までしか追加しません。
PGI_ACC_POOL_ALLOC_MAXSIZE 割り当ての最大サイズを設定します。 割り当てのデフォルトの最大サイズは64Bですが、16B以上であれば別のサイズ(つまり100KB、10MB、250MBなど)を使用できます。
PGI_ACC_POOL_ALLOC_MINSIZE 割り当てブロックの最小サイズを設定します。 デフォルトサイズは16Bですが、他のサイズを使用することもできます。 サイズは16B以上でなければなりません。
PGI_ACC_POOL_THRESHOLD プールアロケータが占めることができるデバイスメモリの合計の割合を設定します。 デフォルトは50%に設定されていますが、他のパーセンテージを使用することもできます。

PGIアクセラレータのコンパイル事例

▶ 使用例1 一般的なGPUアクセラレータ用のコンパイルオプション (PGI 14.1以降のオプション使用)
$ pgf95 -fast -Minfo=accel -acc -ta=tesla f1.f90
main:
     21, Generating copyin(a(1:n))
         Generating copyout(r(1:n))
     22, Loop is parallelizable
         Accelerator kernel generated
         22, !$acc do parallel, vector(256)
					
$ pgcc -fast -Minfo=accel -acc -ta=tesla c1.c
main:
     23, Generating copyin(a[0:n-1])
         Generating copyout(r[0:n-1])
     25, Loop is parallelizable
         Accelerator kernel generated
         25, #pragma acc for parallel, vector(256) 
▶ 使用例2 GPU実行処理時のプロファイルデータを出力
以下は、Linux上での状況を示したものです。"a.out" と言う実行モジュール名は、Linux の
デフォルト名です。Windows 上では、以下の例の場合、デフォルトでは f1.exe、c1.exe と言う名前の
実行バイナリとなります。

$ export PGI_ACC_TIME=1 (PGI13.1以降)
$ pgfortran -fast -Minfo=accel -acc -ta=tesla f1.f90
$ ./a.out
Accelerator Kernel Timing data
  main
    21: region entered 1 time            (GPU計算領域に1回入った)
        time(us): total=75347 init=74604 region=743(GPU初期化時間 init、領域計算部region)
                  kernels=27 data=716 (領域内のカーネル実行時間 27μsec、データ転送716μsec)
        w/o init: total=743 max=743 min=743 avg=743(GPU初期化時間を除いた消費時間)
        22: kernel launched 1 times         (カーネル起動は1回であった)
            grid: [391]  block: [256]              (Block 256のサイズ、グリッドサイズ391)
            time(us): total=27 max=27 min=27 avg=27(カーネルの計算時間)
					
$ export PGI_ACC_TIME=1 (PGI13.1以降)
$ pgcc -fast -Minfo=accel -acc -ta=tesla c1.c
$ ./a.out
Accelerator Kernel Timing data
  main
    23: region entered 1 time
        time(us): total=72900 init=71919 region=981
                  kernels=30 data=951
        w/o init: total=981 max=981 min=981 avg=981
        25: kernel launched 1 times
            grid: [391]  block: [256]
            time(us): total=30 max=30 min=30 avg=30 
▶ 使用例3 GPU付きホストと一般ホスト実行、どちらでも実行可能な PGI Unified Binaryの生成
$ pgfortran -fast -Minfo -acc -ta=tesla,host f1.f90
main:
      1, PGI Unified Binary version for -tp=nehalem-64 -ta=host  (x64ホスト側のコード生成)
     18, Unrolled inner loop 8 times
     22, Generated an alternate loop for the loop
         Generated vector sse code for the loop
         Generated a prefetch instruction for the loop
     26, Generated an alternate loop for the loop
         Generated vector sse code for the loop
         Generated a prefetch instruction for the loop
     30, Loop not vectorized/parallelized: contains call
main:
      1, PGI Unified Binary version for -tp=nehalem-64 -ta=nvidia  (x64+GPU用のコード生成)
     18, Unrolled inner loop 8 times
     21, Generating copyin(a(1:n))
         Generating copyout(r(1:n))
     22, Loop is parallelizable
         Accelerator kernel generated
         22, !$acc do parallel, vector(256)
     26, Generated an alternate loop for the loop
         Generated vector sse code for the loop
         Generated a prefetch instruction for the loop
     30, Loop not vectorized/parallelized: contains call
     
環境変数ACC_DEVICEの値を切り替えることで、PGI Unified Binary のexecutableを
NVIDIA GPUのバイナリを動作させるか、ホスト側バイナリを実行するかを選択できる。
GPUが搭載されているシステムのデフォルトは、ACC_DEVICE=NVIDIA となっている。
GPUが搭載されていないシステムのデフォルトは、ACC_DEVICE=HOST となっている。

●デフォルト実行(GPUを使用)
$ ./a.out (Windows 上では、$ t1.exe)
 Elpased Time (Initialize + Jacobi solver) :      2.841
FORTRAN STOP

●ホストCPU側のコードで実行
$ export ACC_DEVICE=HOST (あるいは host)
$ ./a.out (Windows 上では、$ t1.exe)
 Elpased Time (Initialize + Jacobi solver) :     13.655
FORTRAN STOP

●再度、GPU側で実行       
$ export ACC_DEVICE=NVIDIA (あるいは nvidia)
 Elpased Time (Initialize + Jacobi solver) :      2.841
FORTRAN STOP 
▶ 使用例4 コンパイラが翻訳したGPU用の kernelソースコードのファイルを作成
【プログラム例】
    program main
    use accel_lib
    integer :: n        ! size of the vector
    (中略)
    call system_clock( count=c1 )
    !$acc kernels
        do i = 1,n
            r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2
        enddo
    !$acc end kernels
    call system_clock( count=c2 )
    (中略)
    end
    
    
$ pgfortran -Minfo -acc -ta=tesla,keep f2.f90
main:
     26, Generating copyin(a(1:n))
         Generating copyout(r(1:n))
     27, Loop is parallelizable
         Accelerator kernel generated
         27, !$acc do parallel, vector(256)
         
$ ls
a.out f2.f90 f2.001.gpu (翻訳したkernelのコードのサンプル)

$ cat f2.001.gpu
#include "cuda_runtime.h"
#include "pgi_cuda_runtime.h"
static __constant__ struct{
    int tc4;
    char* p1;
    char* p2;
    }a2;
extern "C" __global__ void
main_27_gpu(
){
float _r_1, _r_2;
int xthreadidx_x;
int xblockidx_x;
int xgriddim_x;
int i1;
int i1s;
xthreadidx_x = threadIdx.x;
xblockidx_x = blockIdx.x;
xgriddim_x = gridDim.x;
i1s = (xblockidx_x)*(256);
if( (i1s) >= (a2.tc4) ) goto _BB_6;
_BB_4: ;
i1 = (xthreadidx_x)+(i1s);
if( (i1) >= (a2.tc4) ) goto _BB_5;
_r_1 = cosf(((float*)a2.p1)[i1]);
_r_2 = sinf(((float*)a2.p1)[i1]);
((float*)a2.p2)[i1] = ((_r_1)*(_r_1))+((_r_2)*(_r_2));
_BB_5: ;
i1s = ((xgriddim_x)*(256))+(i1s);
if( (i1s) < (a2.tc4) ) goto _BB_4;
_BB_6: ;
}

PGIアクセラレータプログラミングモデルの既知の制限事項 (PGI 10.5現在)

  • GPUアクセラレータに処理をオフロード(処理を依頼)するループのネスト(多重ループ)領域は、必ず「rectangular」の形態であること。 特に、triangular ループ、あるいは、以下の例のようなネストした多重ループを成している一方のループの値が、他方のループインデックスの上限・下限値を決めるようなループは、サポートされません。これは、NVIDIA GPU のアーキテクチャに依存した制約です。
    例:

    for (j=0; j<n; J++)
    for (i=0; i<j; i++)
    {some code}

  • GPUアクセラレータに処理をオフロード(処理を依頼)するループの中で配列にアクセスするために使われている「ポインタ」は、C11 言語の 'restrict' 属性を有した宣言を行わなければならない。
     あるいは、safeptrとしたプログラムとして構成するか、あるいは、当該アクセラレータ用のループを含む全体のプログラムファイルに対して -Msafeptr オプションを付加してコンパイルすることでも代替できる。ただし、これらのアプローチは、副作用を伴うため、実行結果の検証を行う必要がある。
  • 少なくとも、オフロードの対象となるループのいくつかは、同期を伴わない、あるいは、イテレーション間の依存性がない完全なデータ並列の特性を有すること。 こうしたループは、NVIDIA GPU 内のマルチプロセッサ間で処理の分割が可能となる。また、ネスト内の一つ以上のループは、同期処理を必要とするようなベクトル処理を行うループとすることができる。例えば、多くのケース、リダクション演算はOKである。こうしたループは、NVIDIA GPU 内の一つのマルチプロセッサの中の複数のプロセッサによりベクトル処理ができる。また、ネスト内の一つ以上のループは、シーケンシャル実行することもできるが、こうしたループは、一つのスレッド・プロセッサ内で順番に実行される。-e.g. 最内側ループのように
  • 配列のインデックスを計算する演算(gatherあるいは scatterのようなもの)は、避けるべきである。ループ・ネスト内にこうした演算がある場合は、ループの並列化あるいはベクトル化を阻止する「依存性」としてコンパイラが認識する。PGI Acceleratorコンパイラの今後のリリースでは、ディレクティブに"independent" 節を導入し、こう言ったループの依存性が存在しないことを明示的に指示できるようにする予定である。この結果、こうしたループも並列化が可能となる。
  • 関数・サブルーチンのコールは、オフロードの対象となるループ内に存在してはならない。いくつかのケースにおいては、コンパイラオプション -Minline を使って、関数をインラインすることもできる。しかしながら、できるだけ、手動で関数等をループ内にインライン展開することで、アクセラレータ領域内で、関数コール文の使用を避けることを推奨します。今後のリリースで、より向上させた自動インライン機能を組み込む予定である。
  • GPUアクセラレータに処理をオフロード(処理を依頼)するループの中では、ポインタ演算はできない