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

 ● CUDA Fortranコンパイルの一例
    pgf95  test.cuf    （最適化なし）
　　pgf95 -O2 test.cuf （最適化あり）
　　pgf95 -O2 -Mfixed test.cuf （最適化あり、ソースはF77固定記述形式）
　　pgf95 -O2 -Mcuda test.f90
 ● エミュレーションモード
　　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ドライバーを実装してあることが必要です）。

  $ pgaccelinfo
  CUDA Driver Version            2030
  （以下省略）
  ここで、
  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以降）と表示されます

• （CUDA Toolkit バージョンのデフォルト値の変更）PGIコンパイラの各リビジョンは、デフォルトで使用する CUDA toolkitのバージョンが予め指定されています。CUDA Fortran あるいは、PGI Acceleratorアプリケーションは、デフォルトでは、この CUDA Toolkit バージョン（上記の各リビジョンのToolkitの説明を参照してください）を利用してビルドします。なお、この使用するCUDA Toolkitのバージョンのデフォルト値の変更が可能です。。コンパイラソフトウェアを実装している $PGI/{target}/{version番号}/bin 配下に siterc ファイルがありますので、その中に、以下を挿入します。ここのパスで言う {target} は、linux86, linux86-64, win32, win64, osx86 等を指します。

  （一例）
  set DEFCUDAVERSION=4.0;
  

  上記の設定を行わず、コンパイラコマンド上のオプションで明示的に 4.0 Toolkit (PGI 11.6 以降）を使用するように指示するには、以下のように指定します。

  pgfortran -ta=nvidia:4.0,.... 
  pgcc -ta=nvidia:4.0,....
  

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

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

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
      iib = (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 :: 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 

　GPU上で動作する PGI CUDA Fortran プログラム内のカーネル・プログラムのデバッグは、一般には行うことができません。これは、ホスト側から GPU 内のdevice memory を直接アクセスできない理由に因るものです。PGIは、仮想的に GPU のカーネルコード部分をホスト側で実行できる「エミュレーション」環境を提供します。この目的は、GPU上の実際の完璧な動作挙動を再現することが目的ではなく、一般には覗き見ることができない「カーネルコード」内の変数や配列等の確認をデバッグ用途で使用できる環境を提供することです。PGDBGは、PGI社の並列デバッガですが、使用法としては、一般のシリアル実行デバッグやマルチコア、MPI等の並列デバッグにおける操作と同じであり、ソースプログラム上に当該ブレークポイントを設定して、対象となる変数や配列の値を確認しながら、デバッグ作業を行う形になります。

【デバッグオプション -g と エミュレートモード -Mcuda=emu を指定】
$ pgf90 -g -Mcuda=emu -o test.exe test.cuf

$ export OMP_NUM_THREADS=4 (仮想的に4スレッドを使用して並列デバッグ）
$ pgdbg test.exe (PGI PGDBG を起動し、デバッグウインドウを開く）

　PGDBGは、マルチスレッド、MPIマルチプロセス対応の並列デバッガです。マルチスレッド・デバッグの場合は、論理的には最大256 スレッド実行までのデバッグが可能ですが、実際は、ハードウェアのCPUコア数に制限があり、実装コア数制限内で並列スレッド・デバッギングを行う事が一般的です（実装コア数より大きな実行スレッドを指定すると、当然ながらデバッガ動作が鈍くなる）。ここでは、CPU上で物理的に動作可能なスレッド数を指定して、「スレッド並列のデバッグ」と同じように実施します。（同時に動作するスレッド並列処理が、後述するいわゆるワープと言う「ベクトル処理」に見立てることもできるように見えます）
　スレッドGPUの挙動は、その基本がスレッド・ブロック内の CUDA Warpsize によるベクトル処理と、スレッド・ブロック群の並行（並列）処理を行う形ですが、ホスト側でこの挙動をそのままエミュレートすることはできません。特に、後者の「スレッド・ブロック」の並列処理は、ホスト側エミュレートでは逐次的処理となりますので、並列動作のエミュレートはできないことに留意して下さい。エミュレーション・デバッグ時には、並列動作挙動が全然異なると言う前提で、デバッグを行う事が必要です。特にデータの依存性に係わる問題のデバッグに関しては、エミュレートモードのデバッグ結果を鵜呑みにするようなことはしないで下さい。

PGDBG を起動して、データ・ウインドウを開いた。

ソースコード内の device subprogram 内にブレークポイントを設定した。

デバイス内の配列数値をデバッグするため表示させた。

ブレークした時点でスコープ可能なローカルデータを表示させる。