PGIアクセラレータ™ コンパイラの使用における FAQ

(2009/12/10記）
PGI 2010 (PGI 10.0) の正式リリース後は、NIVIDIA社の許可により、必要となるCUDAのライブラリは、PGIのインストール時に実装されますので、CUDA Toolkit に関するインストールは行わなくても良いです。最低限、CUDA ドライバーは必ずインストールして下さい。。現在、CUDAソフトウェア（2011年1月現在 CUDA 3.2) は、NVIDIA社のサイトからダウンロードできます。CUDAドライバ環境を構築後、PGIコンパイラをインストールします。
（ご参考）

• NVIDIA CUDA開発環境のインストールに関しての記事をご参照下さい

　この問題は、GNU g++ 4.3 ヘッダーファイルと NVIDIA CUDAソフトウェア間の非互換性によることで生じます。これを改修するには、PGIコンパイラの設定ファイル $PGI/linux86-64/{release}/bin/sitenvrc ファイルに、"set GCCVERSION=40301;" (一例）を追加することで回避できます。もし、sitenvrc ファイルが存在しない場合は、このファイル名で作成して、以下の記述を行って下さい。

まず、GNU g++のバージョンを確かめます。
　kato@photon28:~/GPGPU/PGItest> g++ -v
　(略）
　gcc version 4.3.2 [gcc-4_3-branch revision 141291] (SUSE Linux)
　GCCVERSION が 4.3.2 であることが分かります。これを GCCVERSION=40302; と表記します。

PGI のインストールバージョンが PGI 9.0-4 とした場合、$PGI/linux86-64/9.0-4/bin/sitenvrc ファイルの中に以下の記述を追加記述します。
　set GCCVERSION=40302;

　この問題は、CUDAソフトウェアの libcudart.so ライブラリの互換性がないために生じる問題です。上記の FAQ #001 を参考に、PGIコンパイラの初期設定ファイル siteenvrc の CUDALIBの設定に誤りがないか確認して下さい。64bit Linux の場合、よくある間違いは、set CUDALIB=/usr/local/cuda/lib64;とすべきところを set CUDALIB=/usr/local/cuda/lib; としている場合があります。但し、CUDA 2.2 以前は、CUDAのソフトウェア実装が異なっていたため、64bit Linux の場合でも set CUDALIB=/usr/local/cuda/lib; でした。実際に実装されている CUDA のバージョンをお確かめの上、適切な設定を行って下さい。
　PGI 2010 リリース以降では、cuda ライブラリは、PGIに同梱され、インストール時に実装しますので、こうした問題はないはずです。

　この問題は、NVIDIA CUDAの既知のものであり、CUDAアプリケーションを実行する際に、Linuxカーネル上の CUDA module のロードが行われておらず、そして /dev 上の nvidia エントリが作成されていないために生じます。システムのブート時に明示的に、以下のスクリプトを実行するようにしてください。もし、これにても改善されない場合は NVIDIA CUDA Zone 等に問い合わせて下さい。

#!/bin/bash

# chkconfig: 2345 90 10
### BEGIN INIT INFO
# Provides:       NVIDIA
# Required-Start:
# Required-Stop:
# Default-Start:  3 5
# Default-Stop:
# Description:    CUDA module must be loaded and the entries in /dev created.
### END INIT INFO

modprobe nvidia

if [ "$?" -eq 0 ]; then

# Count the number of NVIDIA controllers found.
N3D=`/sbin/lspci | grep -i NVIDIA | grep "3D controller" | wc -l`
NVGA=`/sbin/lspci | grep -i NVIDIA | grep "VGA compatible controller" | wc -l`

N=`expr $N3D + $NVGA - 1`
for i in `seq 0 $N`; do
 mknod -m 666 /dev/nvidia$i c 195 $i;
done
 mknod -m 666 /dev/nvidiactl c 195 255

else
  exit 1
fi

　ルート権限で、上記の内容を CUDA（一例）と言うファイルに記述します。このファイルを /etc/init.d 配下に置きます。そして、以下の chkconfig コマンドで、Linux ブート時に常に CUDA moduleのロードと /dev/nvidia** の作成を行うように設定します。

[root@photon29 init.d]# pwd
/etc/init.d
[root@photon29 init.d]# ls -l CUDA
-rwxr-xr-x 1 root root 640 12月 22 21:55 CUDA
[root@photon29 init.d]# chkconfig --add CUDA  (boot時にCUDAスクリプトを実行するための設定）
[root@photon29 init.d]# chkconfig --list CUDA
CUDA            0:off   1:off   2:off   3:on    4:off   5:on    6:off

　この問題は、PGI 10.3 Linux版を使用する際、shellのセッションを「日本語のロケール」で環境をセットしている場合に生じます。「日本語のロケール」とは、具体的に言えば、コマンドのメッセージ等で日本語の表示が出るような環境を言います。現在、この問題は弊社から米国 PGI 社に Technical Problem Report として報告していますが、暫定的に以下の方法で、この問題を回避でしてください。

各自のユーザ shell のセッショにおいて、「日本語」ロケールの使用を止め、「英語」ロケールを設定することでこの問題を解決できます。なお、本問題は、PGI 10.4 にて修正されました。

具体的には、$HOME/.bashrc (bash の場合）内に以下を加えて、セッション開始時に環境変数 LANG を
設定する。
　　　　export LANG=C
    

あるいは、コンパイラを使用する時の shell のセッションの開始時に明示的に以下のコマンドを
実行する。
　　　　export LANG=C

(2011/3/11)
　実行できます。
　PGI accelerator コンパイラで GPU 対応のコード生成を行う時は、NVIDIA GPUの世代（Compute Capability：CC と言う) に関するオプションを明示的に指定しない限り、現世代 Fermi から過去の世代までのコードを一つのバイナリの中に作成します。PGI Accerartor Programming model(directive 指定）の場合を例にとり、コンパイルのメッセージで、この状況を説明します。ここでは、HimenoBenchのソースを コンパイルしてみます。-ta=nvidia と言うオプションが、アクセラレータ用の directive を認識し、GPU 用コードを生成せよと言う意味のオプションです。

[kato@photon29 Himeno]$ pgf90 -fastsse -Minfo=accel -ta=nvidia himenoBMTxp.F90
jacobi:
    301, Generating local(wrk2(1:mimax,1:mjmax,1:mkmax))
         Generating copyin(bnd(1:mimax,1:mjmax,1:mkmax))
         Generating copyin(wrk1(1:mimax,1:mjmax,1:mkmax))
         Generating copyin(c(1:mimax,1:mjmax,1:mkmax,1:3))
         Generating copyin(b(1:mimax,1:mjmax,1:mkmax,1:3))
         Generating copyin(a(1:mimax,1:mjmax,1:mkmax,1:4))
         Generating copy(p(1:mimax,1:mjmax,1:mkmax))
    308, Generating compute capability 1.0 binary
         Generating compute capability 1.3 binary
         Generating compute capability 2.0 binary
    310, Loop carried dependence due to exposed use of 'p(2:imax-1,2:jmax-1,2:kmax-1)' 
         Loop carried dependence due to exposed use of 'p(1:imax-1,3:jmax,2:kmax-1)' 
         Loop carried dependence due to exposed use of 'p(3:imax,1:jmax,2:kmax-1)' 
         Loop carried dependence due to exposed use of 'p(1:imax-2,1:jmax-2,2:kmax-1)'
         Loop carried dependence due to exposed use of 'p(2:imax-1,3:jmax,1:kmax-2)' 
         Loop carried dependence due to exposed use of 'p(2:imax-1,2:jmax,3:kmax)' 
         Loop carried dependence due to exposed use of 'p(3:imax,2:jmax-1,3:kmax)' 
         Loop carried dependence due to exposed use of 'p(1:imax-2,2:jmax-1,1:kmax)'
         Loop carried dependence due to exposed use of 'p(2:imax-1,1:jmax-2,1:kmax)' 
         Loop carried dependence due to exposed use of 'p(2:imax,2:jmax-1,1:kmax-2)' 
         Accelerator restriction: scalar variable live-out from loop: gosa
         Parallelization would require privatization of array 'wrk2(2:imax-1,i3+2,i2+2)'
         Sequential loop scheduled on host
    314, Loop is parallelizable
    316, Loop is parallelizable
    318, Loop is parallelizable
         Accelerator kernel generated
        314, !$acc do parallel, vector(2) ! blockidx%y threadidx%z
        316, !$acc do parallel, vector(2) ! blockidx%x threadidx%y
        318, !$acc do vector(128) unroll(2) ! threadidx%x
             Cached references to size [130x4x4] block of 'p'
             CC 1.0 : 32 registers; 8448 shared, 56 constant, 88 local memory bytes; 33% occupancy
             CC 1.3 : 64 registers; 8448 shared, 56 constant, 0 local memory bytes; 25% occupancy
             CC 2.0 : 63 registers; 8328 shared, 140 constant, 0 local memory bytes; 33% occupancy
        332, Sum reduction generated for gosa
    342, Loop is parallelizable
    344, Loop is parallelizable
    346, Loop is parallelizable
         Accelerator kernel generated
        342, !$acc do parallel(16) ! blockidx%y
        344, !$acc do parallel ! blockidx%x
        346, !$acc do vector(128) ! threadidx%x
             CC 1.0 : 13 registers; 72 shared, 48 constant, 0 local memory bytes; 66% occupancy
             CC 1.3 : 13 registers; 72 shared, 48 constant, 0 local memory bytes; 100% occupancy
             CC 2.0 : 18 registers; 8 shared, 88 constant, 0 local memory bytes; 66% occupancy

上記の一番最後の3行に注目してください。

             CC 1.0 : 13 registers; 72 shared, 48 constant, 0 local memory bytes; 66% occupancy
             CC 1.3 : 13 registers; 72 shared, 48 constant, 0 local memory bytes; 100% occupancy
             CC 2.0 : 18 registers; 8 shared, 88 constant, 0 local memory bytes; 66% occupancy

CC 2.0 と言うのは Fermi GPU　で CC 1.3 と言うのは、GTX285系の前世代のCompute Capabilityを意味します。PGIコンパイラは、過去の CC1.0 までのCUDA code をデフォルトで生成し、実行モジュールの中に、生成します。 実行モジュールがGPU搭載マシンで実行する際に、GPUハードウェアの CC を判別し、その CC 用のバイナリコードで実行します。従って、お望みの通りに実行できます。

(2011/3/11)
　上記の #6 の回答で説明したように、PGIコンパイラは、コンパイルするマシンに搭載しているGPUボード専用のコードを生成するのではありません。PGI をインストールしているマシンの搭載ボード特性に依存して実行モジュールを作るのではなく、NVIDIA の GPU の「各世代のバイナリ」を有する Unified Binary が作られるため、実行時には、それぞれの実装された GPU 用のコードが実行されます。こう言った意味で、性能は GPU ネイティブな性能を得ることができます。 但し、こうしたCUDA実行環境を得るためには、以下の注意が必要です。

1. 各マシン共に、CUDA driver のバージョンを合わせること。
2. PGIコンパイラをインストールしていないマシン上には、PGI 専用のshared library(runtime library) をコピーして LD_LIBRARY_PATH 等で、その場所を定義しておくこと（これは Linux 上での基本的事項です）

(2011/4/20)
　NVIDIA GPU のハードウェア仕様で、一つの配列（シングル・オブジェクト）で割付可能なサイズの制約があります。これを Maximum Memory Pitch (memPitch) のサイズと称します。例えば、C2050 や GTX 580 等の Fermi GPUの場合、この値が 2147483647Byte に制約されています。これは、いわゆる、2GB の壁（2**31-1）が存在します。旧世代の GPU の場合は、この値がさらに小さい値で制約されます。
　この技術的仕様に関しては、以下の CUDA プログラミングの仕様書の「デバイス・プロパティ」の memPitch の説明で理解できます。

• memPitch is the maximum pitch in bytes allowed by the memory copy functions that involve memory regions allocated through cudaMallocPitch();
• (日本語）cudaMallocPitch()等の API を通してアロケーションするメモリ領域を含む、メモリコピーを行う関数によって許容される最大のメモリ（ピッチ）のサイズを memPitch と称するが、これに制約がある。）

　Tesla C20XX 系は、memPitchの制限値は、2147483647 bytes です。従って、CUDA C や CUDA Fortran、さらに PGI Accelerator プログラミングモデルを使用する場合、一つの配列の割付の最大サイズは、Maximum Memory Pitch (memPitch)に制約されることにご注意下さい。今の使用している GPU の memPitch の値を知りたい場合は、以下の pgaccelinfo コマンドでデバイスの特性を調べて下さい。

$ pgaccelinfo
CUDA Driver Version:           3020
NVRM version: NVIDIA UNIX x86_64 Kernel Module  260.19.26  Mon Nov 29 00:53:44 PST 2010

Device Number:                 0
Device Name:                   GeForce GTX 580
Device Revision Number:        2.0
Global Memory Size:            1609760768
Number of Multiprocessors:     16
Number of Cores:               512
Concurrent Copy and Execution: Yes
Total Constant Memory:         65536
Total Shared Memory per Block: 49152
Registers per Block:           32768
Warp Size:                     32
Maximum Threads per Block:     1024
Maximum Block Dimensions:      1024, 1024, 64
Maximum Grid Dimensions:       65535 x 65535 x 1
Maximum Memory Pitch:          2147483647B            <====== これです
Texture Alignment:             512B
Clock Rate:                    1544 MHz
Initialization time:           2488 microseconds
Current free memory:           1449000960             <====== 使用可能なfreeメモリ
Upload time (4MB):              898 microseconds ( 719 ms pinned)
Download time:                 1114 microseconds ( 643 ms pinned)
Upload bandwidth:              4670 MB/sec (5833 MB/sec pinned)
Download bandwidth:            3765 MB/sec (6523 MB/sec pinned)

(2011/4/20)
　実際の GPU 上で、どの程度のアロケーションが可能かどうかを確認するための CUDA Fortran で書いた簡単なプログラムを提供します。このプログラムは、実行時の引数に、「アロケーションサイズ(GB)」とそのアロケーションを何回繰り返すかという「カウント数」を指定するとそのサイズを順番にアロケーションして、アロケーションが成功するかどうかを見るためのものです。以下に、一例を示します。この実行例は、GTX 580（1.5GB）上で、0.3GB ずつ、5回アロケーションを繰り返すパターンをテストしたものです。5回目の繰り返しで、フリーメモリが足りないため、エラーとなっています。

このプログラムは PGI 11.4 以上のバージョンでコンパイル・リンクを行う必要があります。
内部で使用している CUDA API が PGI 11.3 以前ではサポートされていません。

(コンパイル）
[kato@photon29 2GB]$ pgfortran mem.cuf

(実行）0.3GB ずつ、5回アロケーションを繰り返す。
[kato@photon29 2GB]$ a.out 0.3 5


                Total Device Memory (B) = 1609760768
              Allocation_unit size (GB) = 0.30
 Real allocation Bytes, Increment count = 322122560    5
 ----------------------------------------------------------------------
 
 Increment No. 1
 Allocation  OK! : Now Free Memory (B) :               1164574720
 ----------------------------------------------------------------------
 Increment No. 2
 Allocation  OK! : Now Free Memory (B) :                842399744
 ----------------------------------------------------------------------
 Increment No. 3
 Allocation  OK! : Now Free Memory (B) :                520224768
 ----------------------------------------------------------------------
 Increment No. 4
 Allocation  OK! : Now Free Memory (B) :                198049792
 ----------------------------------------------------------------------
 Increment No. 5
0: cudaMalloc: 322122560 bytes requested; not enough memory: 2(out of memory)

(プログラム）mem.cuf
program mem_inc
!
! Device Memory allocation test
!    SofTek Systems Inc.  -- Kato
!
  use cudafor
  implicit none
  integer :: i, istat
  integer :: n, nargs
  real (4):: size_gigas
  integer (8) :: size_bytes

  character*20 arg

  type(c_devptr) :: cdvA(5)
  integer(kind=cuda_count_kind) :: total,free

  n = 1
  size_gigas = 2.0

  nargs = command_argument_count()
  if (nargs .ge. 3) then
    print * ,"Usage : a.out {Allocation_size(GB)} {counts}"
    stop
  end if
  do i = 1, nargs
    call get_command_argument(i,arg)
    if (i == 1) read(arg,'(f5.2)') size_gigas
    if (i == 2) read(arg,'(i2)') n
  end do
  if (n .gt. 5) then
    print *, "Inc. count < 6"
    stop
  end if

  size_bytes = size_gigas * 1024 * 1024 * 1024

  print '(/)'
  istat = cudaMemGetInfo(free,total)
  print '(1x,a40,1x,I10)', "Total Device Memory (B) =", total

  print '(1x,a40,f5.2)', "Allocation_unit size (GB) =", size_gigas
  print '(1x,a40,I10,I5)', "Real allocation Bytes, Increment count =", size_bytes, n
  print '(1x,70(1h-)/)'

! Allocations
  do i = 1, n
    print '(a,i2,a)', " Increment No.",i
        istat = cudaMalloc(cdvA(i), size_bytes)
    if (istat == 0) then
          istat = cudaMemGetInfo(free,total)
          print *, "Allocation  OK! : Now Free Memory (B) :", free
        endif
    print '(1x,70(1h-))'
  end do

! Free
  do i = 1, n
    istat = cudafree(cdvA(i))
  end do

  end

(2011/11/11)
　この説明をする前に、Fortran におけるプログラム記述書式について説明が必要です。Fortran では、FORTRAN77 以前の fixed format(実行文は7カラム目から 72カラム目までと言う書式）と Fortran90 以降の Free Format（自由書式）の二つの記述書式があります。古いレガシーなプログラムでは、fixed formatで記述したプログラムが多いと思いますが、この二つの書式の違いによって、アクセラレータ・ディレクティブの「継続」のための書式の書き方が異なります。
　一般に、FORTRAN77 スタイルの fixed format の場合は、その継続行の示し方は 6カラム目に任意の文字を入れる方法となります。!$acc の directive でも同様で、その directive clauses が複数行に渡って定義する場合は、!$acc の継続を行わなければなりません。その方法は以下の通りです。6カラム目に「継続文字」を入れて、継続することを伝えます。ここでは、＋文字を継続文字として使っています。

!$acc region
!$acc+copyin (...) 
!$acc+copyout(...) 

以下のように、7カラム目に空白欄を入れた方が見えやすいかもしれません。
!$acc  region
!$acc+ copyin (...)
!$acc+ copyout(...) 

一方、Fortran90 以降の Free format の記述形式の場合は、一般に、継続する際の文字として "&" を使用します。継続する行の前の行の末尾に "&" を入れることで、継続することを指示します。!$acc の directive でも同様で、末尾に "&" を入れると、次の行は継続することを意味します。

!$acc region &
!$acc copyin (...) &
!$acc copyout(...) 
					

　さて、PGIコンパイラでは、 fixed format のソースファイルであるか、Free format のソースファイルかは、.f（.F) の suffix か、.f90 or .f95 等の suffix かによって区別します。.f の場合だけは、Fixed format となります。必ず、72 カラム形式で記述したプログラムでなければなりません。.f90 (.F90, F95) の場合は、fixed format でも良いし、free format の形式でも良いソースファイルとして認識します。

【参考】 　　PGI コンパイラ 使用ガイド　　P 4 の 1.4 項に記載