PGI 2010のCUDA FortranでのModule文使用時の注意
PGI Accelerator Compiler, CUDA Fortran, Multiple GPUs
Fortran 90 構文でよく利用される MODULE 定義ですが、CUDA Fortran においてその使用方法に注意が必要です。PGI 2010 (PGI 10)のリリース時点では、Module 内に GPU 用の「デバイス変数、デバイス配列」の allocatable 宣言や、あるいは、他のモジュールで定義された GPU 用のデバイス配列の(USE 文による）使用 を行っている Kernel(Global属性) プログラムは、コンパイル時にコンパイラの Internal Error あるいは、その他のエラーを引き起こします。PGI 2010 の時点では、このような使用法は CUDA Fortran の「制約事項」となっております。今後のリリースで対応する予定です。更新 2010/5/11 Copyright © 2010 株式会社ソフテック 加藤
CUDA Fortran コンパイル時に Internal Error となる
PGI 2010 (10.x) の CUDA Fortran では、「独立」の Module 内にデバイス変数・配列データを宣言したグローバル変数・配列を、別の「Fortran Module」で定義された"デバイス・サブプログラム"の中で USE による引用使用することはできない」と言う制約事項があります。この問題の本質的な原因は、NVIDIA社の GPU デバイス・プログラムに対応するリンカーが少なくとも現時点で存在していないことによるものです。具体的に言えば、ある Module で定義された"グローバル変数"を他のモジュールで使用するために「そのシンボル」を結合できないと言うことに因るものです。（一般に、商用コンパイラのリンカー部分は、Linux or OS X では GNU のリンカーやWindows であればMicrosoftの SDK のリンカーを使用しています。）この制約に関しては、2010 年 5 月時点において、PGI によりその開発が進められておりますが、PGI の今後のリビジョン（完全にサポートするのは PGI 11.0 以降で行います。）で対応する予定です。（後記：この中の一部の制約事項をなくすために、 10.6 リリースで機能追加しました。）
こうしたプログラムを CUDA Fortran (pgf90） でコンパイルすると、Internal Error等のメッセージエラーメッセージで、問題となるエラー箇所が指摘されない状況となっています（この点に関しては、現在、Technical Problem Report として issue しています）。今後、CUDA Fortran 構文チェックにおけるメッセージも改善されると思われますが、以下のような原因を特定できないエラーが生じた場合は、再度、CUDA Fortran の構文や制約事項の確認をお願いします。
PGI CUDA Fortran Programming Guideによる現在の仕様（PGI 10.x 上の機能) は、「デバイス・サブプログラム（いわゆる attribute(global) , attribute(divice) 属性プログラム）は、他の Module で宣言された"デバイス変数"のアクセスにはできず、「同じモジュール」内で宣言された"デバイス変数"あるいは、"形状固定のデバイス配列(サイズが固定したもの）"しかアクセスできない」と言う制約があります。従って、allocatable な"デバイス配列"等を使用する一般的なプログラムの書き方としては、ホスト側領域から「引数」でデバイス・サブプログラムにデータを渡す形となります（以下の例）。但し、その引数文字数の長さは 256byte と言う制約があることも忘れてはなりません。
（一例） PGF90-F-0000-Internal compiler error. pgnvd job exited with nonzero status code 0(xxx.cuf:30) PGF90-F-0000-Internal compiler error. unexpected runtime function call 0(yyy.CUF: 64)
以下に、PGI 10.x の全てのバージョンでエラーとなる、モジュール内に「デバイス変数」を定義した場合の例を示します。これは、この独立したモジュールを別のデバイス・サブプログラム内で USE 文の指定を行い、使用する形態となっています。
【エラーとなるプログラム例】 test0.cuf
! モジュール data の中で int_d と言う動的割付型デバイス配列を宣言
module data
! PGI 10.x CUDA Fortran don't allow "device,allocatable"
! this is a limitation in current ... SofTek
integer, device, allocatable, dimension(:) :: int_d
end module data
! モジュール cudamod は、kernel ルーチンを含むモジュール
module cudamod
use cudafor
use data ! モジュール内で device 属性の変数は使えない
implicit none
contains
! kernel subprogram
attributes(global) subroutine test1 (n)
integer :: it, ib
integer, value :: n
!! integer, device :: int_d(n)
it = threadidx%x
ib = (blockidx%x-1) * 16
int_d(it+ib) = it + ib
end subroutine test1
end module cudamod
program cuda_device
use cudafor
use data
use cudamod
implicit none
integer :: n=64
integer :: int_h(n), ist
int_h = 0
allocate(int_d(n))
print *, "n/16",n/16
call test1 <<<n/16,16>>> (n)
ist = cudathreadsynchronize()
int_h = int_d
print *,'int_h = ',int_h
deallocate(int_d)
end program cuda_device
--------------------------------------------------------------------------------------
[kato@photon29 IN-MODULE]$ pgf90 test0.cuf
/tmp/pgcudaforqHVcW1IGl2O9.gpu(4): error: variable "data_16" has already been defined
1 error detected in the compilation of "/tmp/pgnvdAIVcoZyoiwqX.nv0".
PGF90-F-0000-Internal compiler error. pgnvd job exited with nonzero status code 0 (test0.cuf: 22)
PGF90/x86-64 Linux 10.1-0: compilation aborted
上記のプログラムを以下の test1.cuf あるいは、 test2.cuf のように変更すると、コンパイルが可能となります。すなわち、test1.cuf では、kernel サブプログラムを定義したモジュール内では、デバイス変数を含んだデータモジュールを USE で参照しないようにし、kernel にはデバイス変数を引数で渡すと言う方法です。 これは、一般的な方法で、デバイス変数は親ルーチンで宣言し、アロケートし、kernelへは引数で渡すようにします(test2.cuf)。
test1.cuf は PGI 10.1 までは、実行できましたが本来の仕様から外れたものでした。正確には、test2.cuf のスタイルが本来の仕様です。（2010/3/24追記)
【プログラムファイル】test1.cuf PGI 10.x では、このプログラムは動作しません。 module data !このモジュール内に device 配列が存在します。現在の仕様外の使い方です。 integer, device, allocatable, dimension(:) :: int_d end module data module cudamod use cudafor implicit none contains ! kernel subprogram ! indat 引数でデータを渡す 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 data use cudamod implicit none integer :: n=64 integer :: int_h(n), ist int_h = 0 allocate(int_d(n)) print *, "n/16",n/16 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
【プログラムファイル】test2.cuf
module cudamod
use cudafor
contains
! kernel subprogram
attributes(global) subroutine test1 (intdat,n)
implicit none
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 :: n=64
integer :: int_h(n), ist
integer, device, allocatable, dimension(:) :: int_d
int_h = 0
allocate(int_d(n))
print *, "n/16",n/16
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
[kato@photon29 IN-MODULE]$ pgf90 test2.cuf
[kato@photon29 IN-MODULE]$ a.out
n/16 4
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
