PGI CUDA Fortran のコンパイル・オプション

　PGI CUDA Fortran 構文を含むプログラムのファイル名は、***.cuf と言う名称にします。コンパイラは、この名称のファイルを CUDA Fortran であると認識し、-Mcuda オプションを付けなくてもコンパイルできます。また、***.CUF と cuf サフィックスを大文字とすると、コンパイラは、CUDA Fortran プログラムで、かつ、プリプロセッシング処理を行うべきプログラムであると認識します。もちろん、Fortran のファイル名のサフィックスを *.f、*.F、*.f90、*.F90 と言った従来の慣習である名称としても良いですが、この場合は、コマンドラインに必ず、オプション -Mcuda を明示的に記述しなければなりません。

● CUDA Fortran ファイル名一例
  　*.cuf、*.CUF（CUDA Fortran プログラムファイルであることを明示）
  　*.f、*.F、*.f90、*.F90、*.f95、*.F95 （従来の慣習名）

　PGI pgfortran (pgf95、あるいは、pgf90は同じコンパイラです）コンパイラを使用して、CUDA Fortran プログラムをコンパイルするためのオプションの例を示します。なお、CUDA Fortran は、デフォルトでは F90 以降の「自由記述形式」のファイルと見なしてコンパイルしますので、もし旧 F77 時代の「固定記述形式（7カラムから実行文と言った形式）」の場合は、必ず、-Mfixed と言うオプションを付けて、コンパイラに指示する必要があります。以下は、pgf95（pgfortran、あるいは、pgf90 も同じコンパイラです）を使用した場合の例です。

　PGI 17.1 以降、Pascal(CC 6.0) GPU 用と過去の CC 2.0 用の executable はデフォルトで生成されません。これらの GPU 用の executable の生成には、-Mcuda=cc60 or -Mcuda=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用のコードを生成できます。

● CUDA Fortranコンパイルの一例
    pgfortran  test.cuf    （最適化なし）
　　pgfortran -O2 test.cuf （最適化あり）
　　pgfortran -O2 -Mfixed test.cuf （最適化あり、ソースはF77固定記述形式）
　　pgfortran -O2 -Mcuda test.f90
　　pgfortran -O2 -Mcuda=cuda10.0,cc70 test.cuf (CUDA10.0 toolkit使用, cc 7.0用）
 ● エミュレーションモード (PGI 18.1から機能廃止）
　　pgf95 -O2 -Mcuda=emu test.cuf

• -Mcuda　：　このオプションを指定すると、コンパイラは、一般的な Fortran 構文だけでなく CUDA Fortran 構文を解釈するコンパイラモードとなります。CUDA Fortran プログラムをコンパイルし、必要なライブラリをリンクします。なお、リンク時においてもこのオプションが必要です。
• -Mcuda[=以下のサブオプション] があります。このサブオプションは、カンマ（,）で区切って複数のものを指定することができます。

• （注意） -Mcuda=emu によるエミュレーションモードでの実行は、実際の GPU 上で実行する状況を完全に再現できないことに注意して下さい。特に、エミュレーションモードでは、一時に「シングル」スレッド・ブロックしか実行できません。この状態では、メモリ・アクセス（転送）時に生じるメモリレース等によるエラーは再現できません。また、ホストCPU側の浮動小数点演算器と内部組込関数が使用されるため、GPUデバイス上の演算器等を使用した結果に較べ、若干の数値的差異が起こりえます。
• （更新情報：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.x (2011)では、デフォルトで 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を使用したい場合は、明示的に cuda8.0 or cuda9.0 のサブオプションを指定してコンパイル・リンクする必要があります（但し、CUDA 9.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を使用したい場合は、明示的に cuda8.0 or cuda9.0 or cuad9.2 のサブオプションを指定してコンパイル・リンクする必要があります（但し、CUDA 9.0 or 9.1 oe 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 -Mcuda=cuda9.2,.... 
  pgcc -Mcuda=cuda9.2,....

• PGI の各リビジョンで追加・変更されたオプションについては、最新の「リリースノート」をご覧下さい。

以下に、CUDA Fortran の簡単なプログラム例とコンパイル・実行例を示します。以下のプログラムは、デバイス上の配列(int_d)を宣言して、GPU上で整数配列に値を入れ、ホスト側の配列に戻して印字すると言う単純なものです。しかし、CUDA Fortran 構文の基本的なものが使用されている例です。

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

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

set DEFCOMPUTECAP=60 70;

▶ 簡単なプログラム例

module cudamod
   use cudafor
   implicit none

   contains
   !  kernel subprogram
   attributes(global) subroutine test1 (intdat,n)
      integer :: it, ib
      integer, value :: n
      integer, device :: intdat(n)
      !----------------------------
      it = threadidx%x
      ib = (blockidx%x-1) * 16
      intdat(it+ib) = it + ib
      !----------------------------
   end subroutine test1
end module cudamod

program cuda_device
   use cudafor
   use cudamod
   implicit none
   integer, parameter :: n=64
   integer :: int_h(n), ist
   ! デバイス側の配列データを allocatable で宣言
   integer, device, allocatable, dimension(:)  :: int_d

   int_h = 0
   allocate(int_d(n))
   
   ! カーネル（デバイスプログラム）を起動する
   call test1 <<<n/16,16>>> (int_d, n) !pass arguments
   ist = cudathreadsynchronize()  ! 同期ポイント

   ! デバイス側の配列データをホスト側に戻す
   int_h = int_d
   print *,'int_h = ',int_h
   deallocate(int_d)
end program cuda_device

$ pgf90 -O2 -Minfo -Mcuda -o test.exe test.cuf
cuda_device:
     28, Memory zero idiom, loop replaced by call to __c_mzero4
$ test.exe
 int_h =             1            2            3            4            5
            6            7            8            9           10           11
           12           13           14           15           16           17
           18           19           20           21           22           23
           24           25           26           27           28           29
           30           31           32           33           34           35
           36           37           38           39           40           41
           42           43           44           45           46           47
           48           49           50           51           52           53
           54           55           56           57           58           59
           60           61           62           63           64