PGIアクセラレータ・コンパイル用のオプション

　PGI の F95, C のコンパイラを使用して、PGIアクセラレータ用のコンパイルを行うオプションの例を示します。以下は、pgf95（pgfortran、あるいは、pgf90 も同じコンパイラです）を使用した場合の例ですが、C言語用の pgcc コンパイラのオプションの設定方法も、同様です。なお、コマンドライン上でリンク時にも、必ず、コンパイル時に指定したものと同じ -ta オプションを指定することが必要です。

 ● Fortranコンパイルの一例
　　pgf95 -fast -Minfo -ta=navidia test.f
　　pgf95 -fast -Minfo=accel -ta=nvidia:keepgpu,time test.f90
 ● C99 コンパイルの一例
　　pgcc -fast -Minfo -ta=navidia test.c
　　pgcc -fast -Minfo=accel -ta=nvidia:keepgpu,nofma,time test.c 

• -Minfo=accel　：　このオプションを指定すると、コンパイラがアクセラレータ領域をGPUカーネルに翻訳できたかどうかについて、コンパイラのメッセージとして出力します。-Minfo のみの指定では、その他の最適化情報も併せてメッセージとして出力します。
• -ta=nvidia(,nvidia_suboptions), host　： PGI アクセラレータ・コンパイラに伴う新しいオプションです。-ta は、ターゲット・アーキテクチャを意味します。現時点では、「nvidia」 のみとなります。Fortranにおける !$ACC ディレクティブ、C における #pragma acc ディレクティブをコンパイラに認識させるために、このオプションを使用します。これは、Fortran 並びに C コンパイラのみで有効です。このオプションは、以下のサブオプションを有します。
• -ta=nvidia - NVIDIAアクセラレータをターゲットとして選択します。さらに、以下の nvidia 用のサブオプションがあります。このサブオプションは、カンマ（,）で区切って複数のものを指定することができます。

• -ta=nvidia,host - ターゲットとして、hostを選択する。nvidiaオプションとの組み合わせで使用されます。アクセラレータ領域をホスト側で実行するようにコンパイルする。このオプションは、GPUが実装されていないシステムでも動作するような実行バイナリとなるPGI Unified Binaryコードを生成します。この host サブオプションは、上記の nvidia 専用のサブオプションを指定した後に指定します（一番最後に指定します）。
• （更新情報：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ドライバーを実装してあることが必要です）。

  $ 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 ソフトウェアのツールを発動し、GPU カーネルコードを生成し、オブジェクト・ファイルの中にカーネルを埋め込みます。

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

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

  export ACC_DEVICE=NVIDIA
  setenv ACC_DEVICE NVIDIA

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

  export ACC_DEVICE_NUM=1
  setenv ACC_DEVICE_NUM 1

• ACC_NOTIFY - ACC_NOTIFY 環境変数は、Kernel がアクセラレータ上で実行された際に、そのイベントを標準出力としてショートメッセージで印字するために使用されます。環境変数の値は、負の整数であってはなりません。0を指定するとこの機能を抑止します（デフォルト）。0以外の正数の場合は、カーネルを実行する毎に、ショートメッセージを標準出力に印字します。
  例：
  

  export ACC_NOTIFY=1
  setenv ACC_NOTIFY 1
  (メッセージ一例）
  launch kernel file=/home/kato/GPGPU/OpenMP/jacobi4.F function=jacobi line=229 device=1 grid=2500 block=128x4

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

$ pgf95 -fast -Minfo=accel -ta=nvidia 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 -ta=nvidia 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 と言う名前の
実行バイナリとなります。

$ pgf95 -fast -Minfo=accel -ta=nvidia,time 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（カーネルの計算時間）
					
$ pgcc -fast -Minfo=accel -ta=nvidia,time 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 

$ pgf95 -fast -Minfo -ta=nvidia,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 region
        do i = 1,n
            r(i) = sin(a(i)) ** 2 + cos(a(i)) ** 2
        enddo
    !$acc end region
    call system_clock( count=c2 )
    （中略）
    end
    
    
$ pgf95 -Minfo -ta=nvidia,keepgpu 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: ;
}