PGI 2011のCUDA FortranでのModule文
PGI Accelerator Compiler, CUDA Fortran, Multiple GPUs
CUDA Fortran 上の Module 文の使用自由度が向上しました。PGI 2010 (PGI 10.x)では、Module 内に定義された GPU 用の「デバイス変数、デバイス配列」を使用できる範囲が限定されておりました。「デバイス変数、デバイス配列」は、任意のホスト側のプログラムで使用（USE文の指定は必要）できますが、デバイス側のサブプログラムでは、「同じ Module」で定義された global あるいは、device サブプログラムしか使用できないと言う制約がありました。これに関する詳細は、過去のコラムを参考にして下さい。 PGI 2011 (11.0) 以降では、同じ Moudle 内に無くとも、他の Module ブロック内に定義された global あるいは、device サブプログラムでも、 use 文を指定してアクセスできるようになりました。2011/5/9 Copyright © 2011 株式会社ソフテック 加藤
CUDA Fortran での Module 内の"デバイス変数"の扱い
PGI 2010 (10.0～10.9) の CUDA Fortran では、「独立」の Module 内にデバイス変数・配列データを宣言したグローバル変数・配列を、「別の」Module内に定義された"デバイス・サブプログラム"から USE文 を使って引用使用することができない」と言う制約事項がありました。
PGI 2011（PGI 11.0 以降）では、「別の MODULE 内に定義されている"デバイス変数(配列）"は、このモジュールデータを USE 文を使用することで、任意の global あるいは、device サブプログラム内でシームレスに参照できる」ようになりました。これは、コンパイラが、こうしたモジュールデータを動的にサブプログラムにリンクできるようにコード生成できるようにしたため可能となりました。この動的リンクのタイミングは、最初にデバイス・サブプログラムが起動された時に行われます。但し、この機能は、Fermiアーキテクチャ・クラスの GPU に対して有効ですので、ご注意下さい（Fermi 以降の indirection 機能を使用しています）。
なお、次の技術的ステップとしては、Module 内に定義された「デバイス・サブプログラム」を他のモジュール内の global or device サブプログラムから call できるようにすることですが、これについてはかなり難しい技術的な課題が存在します（CUDAアーキテクチャ、ソフトウェア環境上でのフレキシブルなポインタの扱いが難しい問題）。これに関しても何らかの方法を見出す予定ですが、PGI 2011 リリースにおいては、その機能を提供できません。
以下のプログラムは、モジュール内に定義されたint_d 配列を異なるモジュール内で使用する形の例を示します。以下の構成のプログラムは、PGI 2010(PGI 10.x) では、実行できませんでした。
【プログラムファイル】test1.cuf
module data
integer, device, allocatable, dimension(:) :: int_d
end module data
module cudamod
use cudafor
use data
implicit none
contains
! kernel subprogram
attributes(global) subroutine test1 (n)
integer :: it, ib
integer, value :: 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, parameter :: n=64
integer :: int_h(n)
int_h = 0
allocate(int_d(n))
print *, "n/16",n/16
call test1<<<n/16,16>>> (n)
int_h = int_d
print *,'int_h = ',int_h
deallocate(int_d)
end program cuda_device
●PGI 11.0 以降のコンパイラでコンパイルして実行
[kato@photon29 IN-MODULE]$ pgfortran test1.cuf -V11.0
[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
●PGI 10.9 以前のコンパイラでコンパイルして実行(エラー）
[kato@photon29 IN-MODULE]$ pgfortran test1.cuf -V10.9
[kato@photon29 IN-MODULE]$ a.out
n/16 4
0: copyout Memcpy (host=0xa624210, dev=0x5200000, size=256) FAILED: 4(unspecified launch failure)
