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) 用のサブオプションがあります。このサブオプションは、カンマ（,）で区切って複数のものを指定することができます。

• -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以降）

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

set DEFCOMPUTECAP=60 70;

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アクセラレータのコンパイル事例

$ 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) 

以下は、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 

$ 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 

【プログラム例】
    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: ;
}