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

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

対象 PGI CUDA Fortran

サマリー

CUDA™ は NVIDIA 社の GPU のアーキテクチャですが、NVIDIA 社からのCUDA開発環境は、CUDA C として知られている拡張 C コンパイラとツール群のみが提供されております。CUDA C は、高級言語上から GPU のために CUDA API を使用して明示的にプログラミングすることができるものです。PGI とNVIDIA社は共同で PGI CUDA Fortran の開発を行い、CUDA C と同等な機能を PGI Fortran コンパイラに実装しました。ここでは、PGI CUDA Fortran のコンパイル方法を説明します。PGI 18.7 用更新 2018/8/21Copyright © 株式会社ソフテック

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=cuda9.2,cc70 test.cuf (CUDA9.2 toolkit使用, cc 7.0用)
 ● エミュレーションモード (PGI 18.1から機能廃止)
  pgf95 -O2 -Mcuda=emu test.cuf
Pascal (CC6.0)/Volta(cc7.0) 用の executable を作成する場合は、CUDA 8.0 or CUDA 9.0 以降のライブラリを必要としますので、必ず、-ta=tesla,cuda8.0 or -ta=tesla,cuda9.0 or -ta=tesla,cuda9.2 あるいは、-ta=tesla,cc60 or -ta=tesla,cc70 のオプションを明示的に付けてコンパイルしてください。
  • -Mcuda : このオプションを指定すると、コンパイラは、一般的な Fortran 構文だけでなく CUDA Fortran 構文を解釈するコンパイラモードとなります。CUDA Fortran プログラムをコンパイルし、必要なライブラリをリンクします。なお、リンク時においてもこのオプションが必要です。
  • -Mcuda[=以下のサブオプション] があります。このサブオプションは、カンマ(,)で区切って複数のものを指定することができます。
サブオプション nvidia用 機能
emu エミュレーションモードでコンパイルします。これは、GPU 用のコード生成は行わず、ホスト側でエミュレーション実行可能なコードを生成します。一般に、デバッグ時に使用します。CUDA Fortran の " device code (kernel)" は、ホスト上で実行出来るコードで生成され、ホスト側の pgdbg デバッガを使用できます。(PGI 18.1以降廃止)
cc10 compute capability 1.0 のコードを生成 (PGI 13.10以降廃止)
cc11 compute capability 1.1 のコードを生成 (PGI 13.10以降廃止)
cc12 compute capability 1.2 のコードを生成 (PGI 13.10以降廃止)
cc13 compute capability 1.3 のコードを生成 (PGI 13.10以降廃止)
cc1x compute capability 1.x のコードを生成 (PGI 15.1以降廃止)
cc1+ compute capability 1.x, 2.x, 3.x のコードを生成 (PGI 14.1以降) (PGI 15.1以降廃止)
tesla compute capability 1.x (=cc1x) のコードを生成 (PGI 13.1以降) (PGI 15.1以降廃止)
cc20 compute capability 2.0 のコードを生成 (PGI 10.4以降) (PGI 13.10以降廃止) (PGI 15.5以降復活)
cc2x compute capability 2.x のコードを生成 (PGI 10.4以降)
cc2+ compute capability 2.x, 3.x のコードを生成 (PGI 14.1以降)
charstring GPUカーネル内で文字列の使用を制限付きで使用する(PGI 15.1以降)
felmi compute capability 2.x (=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以降)
kepler compute capability 3.x (=cc3x) のコードを生成 (PGI 13.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以降)
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.1 or 3.1 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以降)
fastmath fast mathライブラリを使用 (PGI 10.4以降)
[no]flushz GPU上の浮動小数点演算の flush-to-zero モードを制御。デフォルトはnoflushz。 (PGI 11.5以降)
keepbin kernelバイナリファイルを保持し、ファイル(.bin)として出力する
keepgpu kernelソースファイルを保持し、ファイル(.gpu)として出力する (PGI 10.3新設)
keepptx GPUコードのためのportable assembly(.ptx)ファイルを保持し、ファイルとして出力する
[no]lineinfo GPU line informationを生成する(PGI 15.1以降、Linux only)
[no]llvm llvm ベースのバックエンドを使用してコードを生成する。(PGI 15.1以降) 64-bit上ではLLVMバックエンドを使う [使わない]。なお、PGI 15.1 以降はデフォルト llvm を使うように変更された。リンク時にエラーが生じる場合、nollvm を試すことをお勧めする。
maxregcount:n GPU上で使用するレジスタの最大数を指定。ブランクの場合は、制約が無いと解釈する
nofma fused-multiply-add命令を生成しない (PGI 10.4以降)
noL1
noL1cache
グローバル変数をキャッシュするためのハードウェア L1 データキャッシュの使用を抑止する (PGI 13.10以降)
loadcache:L1
loadcache:L2
グローバル変数をキャッシュするためにハードウェア L1 あるいは L2データキャッシュを使用する。但し、アーキテクチャ上、有効とならない GPU がある (PGI 14.4以降)
ptxinfo コンパイル時にPTXAS情報メッセージを表示する(PGI 11.0以降)
[no]rdc Fortran Module 内の device routine など、異なるファイルに配置されたデバイスルーチンをそれぞれ分割コンパイルし、リンクが出来るようにする。CUDA 5.0 以降の機能を使用します。(PGI 13.1以降 + CUDA 5.0 以降)(PGI 14.1 以降デフォルト)
[no]unroll 自動的に最内側ループのアンローリングを行う (default at -O3)(PGI 14.9以降)
  • (注意) -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のバージョン制御をご覧ください。
    $ pgaccelinfo
    CUDA Driver Version            9020
    (以下省略)
    ここで、
    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以降)と表示されます
  • (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以降)

 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 定義を追加できます。デバイスコードの生成には時間がかかることがあるため、コンピューティング機能の数が増えるとコンパイル時間が長くなることがあります。

▶ 簡単なプログラム例

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