PGIアクセラレータコンパイラ > OpenACCサイト・トップ > インデックス

OpenACC ディレクティブによるプログラミング

10-2章　OpenACC 2.0 declare data ディレクティブによるハンドリング

declare data ディレクティブによるグローバルデータのハンドリング

　PGI 14.7 から、C のグローバル変数、Fortran Module 変数を対象とした OpenACC declare ディレクティブがサポートされた。以下の図に示す機能を提供する。これは、異なるソースファイルに記述されたルーチン間に跨るグローバル変数を管理するために必須となる機能となる。　グローバルデータのハンドリングでは、copyin、create、device_resident、link、deviceptr の clause が使用出来る。主な clause の使い分けに関して、以下の図でその概略を説明する。なお、declare ディレクティブ自体の説明は、9章を参照のこと。

C のグローバル変数、Fortran Module 変数用の declare data ディレクティブ

　C のグローバル変数、Fortran Module 変数を対象とした OpenACC declare ディレクティブに関して解説する。この用途における declare ディレクティブは、copyin、create、device_resident、link、deviceptr の各 clause と共に使用することができる。なお、これらは NVIDIA tesla GPUターゲットに対して有効である（AMDターゲットは有効ではない）。この機能は、基本的に OpenACC routine ディレクティブや分離コンパイル時に利用されるものであるが、declare clause 内のデータは、プログラムがデバイスにアタッチされる時点で、デバイス上に「静的に」アロケートされる。copyin clause で指定されたデータはその時点で、ホスト側データからコピーされ初期化される。なお、一般にプログラムがデバイスにアタッチするタイミングは、最初に当該データに到達する、あるいは compute 構文に到達した時点、あるいは、OpenACC acc_init API ルーチンを call した時点となる。

　以下に示す C プログラムで説明しよう。以下のプログラムでは、グローバル struct とグローバル配列ポインタを使用している。

struct{
	float a, b;
}coef;
float* x;        // x のポインタ
#pragma acc declare create(coef,x)　//　グローバル変数を定義した後に declare 宣言、割付
. . .
#pragma acc routine seq
void modxi( int i ){
	x[i] *= coef.a;
}
. . .
void initcoef( float a, float b ){
	coef.a = a;
	coef.b = b;
#pragma acc update device(coef)  // デバイス側へ coef データを更新（データコピー）
}
. . .
void allocx( int n ){
	x = (float*)malloc( sizeof(float)*n ); // x の実体の割付
	#pragma acc enter data create(x[0:n]) // x 実体データをデバイス側で割付、データ領域開始
}
. . .
void modx( int s, int e ){
	#pragma acc parallel loop
	for( int i = s; i < e; ++i ) modxi(i);
}

　declare create(coef,x) は、デバイス上に静的に、coef と x ポインタのコピーを割付ける。 initcoef ルーチンでは、ホスト上で coef データが定義され、update directive を指定してデバイス側にそのコピーを行う。allocx ルーチンは、ホスト上の x ベクトルデータのためのスペースを割付し、また acc enter data create(x[0:n]) を使用して、同様にデバイス上に割付けている。x のポインタは、すでに静的にデバイス上で存在しているため、これによって x のデバイスへのコピーはポインタも含めてデバイスデータの更新が行われる。最後に、parallel ループは、グローバル x ポインタと coef struct を参照する modxi ルーチンをコールしている。ホスト上で call された時、このルーチンはホスト上のグローバル x と coef を参照することになる。また、この parallel ループのようにデバイス上で call された場合は、デバイス上のグローバル x と coef を参照することになる。

　もし、modxi ルーチンが別のファイル内に存在した場合、coef と x の宣言は、 external 属性を持つことになる。その場合も、コード自体は下記のように同じように記述する。なお、declare create ディレクティブは、当該変数が extern で宣言されていたとしても、このソースファイルの中では必ず明記する必要がある。

extern struct{
	float a, b;
}coef;
extern float* x;
#pragma acc declare create(coef,x)
. . .
#pragma acc routine seq
void modxi( int i ){
	x[i] *= coef.a;
}

　グローバル変数がデバイスメモリ内に存在すると、それは、ホストとデバイスオブジェクト間の対応表である OpenACC ランタイム "present" テーブル中にも記録されている。こうした理由で、グローバル変数へのポインタは別のファイルにあるルーチンへ引数として渡すことができる。その引数は、あたかも present clause で指定されたポインタとして使うことができる。次は、コールする側のルーチンで、あるサイズの静的なグローバルな係数配列 xcoef を使った例である。

float xcoef[11] = { 1.0, 2.0, 1.5, 3.5, ... 9.0 };
#pragma acc declare copyin(xcoef)
. . .
extern void test( float*, float*, float*, n );
. . .
void caller( float* x, float* y, int n ){
	#pragma acc data copy( x[0:n], y[0:n] )
	{
		. . .
		test( x, y, xcoef, n );
		. . .
	}
}

　上記の #pragma acc declare copyin ディレクティブは、コンパイラに対して、プログラムがデバイスをアタッチする時にホスト側の配列からデバイス配列にコピーして初期化するコードを生成することを指示するものである。一方、以下のように別ファイルに test と言う procedure が記述されており、そのルーチンへの引数の全ては、すでにデバイス上に present(存在）していると言う状態を指示している。この理由は、x と y に関してはすでに caller 側で data 構文を使ってcopyin されており、グローバルな xcoef 配列に関しては、静的にデバイス上に copyin されているからである。

void test( float* xx, float* yy, float* cc, int n ){
	#pragma acc data present( xx[0:n], y[00:n], cc[0:11] )
	{
		. . .
		#pragma acc parallel loop
		for( int i = 5; i < n-5; ++i ){
			float t = 0.0;
			for( int j = -5; j <= 5; ++j ){
				t += cc[j+5]*yy[i+j];
			}
			xx[i] /= t;
		}
		. . .
	}
}

　Fortran においては、MODULE 内の固定サイズ変数、配列と、MODULE のスコープ範囲内で declare ディレクティブ内に指定された MODULE アロケータブル配列が、グローバル変数として CPU 側だけでなくデバイスコード内で使用出来る。declare create, declare copyin あるいは declare device_resident 内で現れる MODULE アロケータブル配列は、allocate 文内にそれらが現れた時点で、デバイスメモリだけではなくホスト側のメモリ内に割り付けられる。コンパイラは、データへの実際のポインタと配列の各添え字の上下限値を含むデスクリプタを管理する。そして、ポインタのデバイスコピーを行うとデバイスメモリ内の配列にひも付けするようにセットされる。

　次に示す Module の例は、ある固定サイズの配列とアロケータブルな配列を含んだもので、両方とも declare create clause で指定している。この場合、静的配列 xstat は、アクセラレータ compute 領域あるいは routines の内部でいつでも使える変数となる。

module staticmod
integer, parameter :: maxl = 100000
real, dimension(maxl) :: xstat
real, dimension(:), allocatable :: yalloc
!$acc declare create(xstat,yalloc)
end module

　当然、この Module は、yalloc 配列をアロケートする「別のファイル」の中で使用される場合もある。実際の割付のタイミングは yalloc がアロケートされる時であり、ホスト側とデバイス側の両方で割付が起こる。そして、それ以降、アクセラレータ compute 領域あるいは routines の内で任意に使えるようになる。

subroutine allocit(n)
use staticmod
integer :: n
allocate( yalloc(n) )
end subroutine

　以下の例は、これらの配列がアクセラレータ compute 領域あるいは routines の内で使用される例である。

module useit
 use staticmod
contains
 subroutine computer( n )
  integer :: n
  integer :: i
  !$acc parallel loop
   do i = 1, n
     yalloc(i) = iprocess( i )
   enddo
 end subroutine
 real function iprocess( i )
  !$acc routine seq
  integer :: i
  iprocess = yalloc(i) + 2*xstat(i)
 end function
end module

declareディレクティブの制限事項

変数、配列は、function, subroutine, program, or module に対する declare directive の全ての clause の中で少なくとも一度指定する。
サブ配列の指定はできない。
Fortran では、大きさ引継ぎ配列(assumed-size array)の仮引数は、declareディレクティブに現れてはならない。
Fortran では、ポインタ配列は指定しても良いが、ポインタ結合（pointer association）は、デバイスメモリ内に保持できない。
Fortran module 宣言部では、create, copyin, deviceptr, device_resident とlink の clause だけが許される。
C/C++のグローバル・スコープ下では、create, copyin, deviceptr, device_resident とlink の clause だけが許される
C/C++の extern 変数は、create, copyin, deviceptr, device_resident とlink の clause 内で指定できる。

Fortran Module グローバル変数を使用する例

　以下の例は、Fortran MODULE 変数を別ファイルで記述されたプログラムで使用する例である。MODULE global_data において、a, b, c というグローバル配列が定義されている。この変数を明示的に declare create ディレクティブで、ホスト側とデバイス上に割付を行う指示をする。これによって、グローバル変数の配列スペースの確保が行われる。次に、declare_test メインプログラムでは、ホスト側で a, b 配列値の定義が行われた後、update device (a,b) ディレクティブを使って明示的にデバイス側へデータのコピー(更新）を行う。なお、update ディレクティブは「実行文」としての機能を果たす。その後の Parallel 構文領域内のループで call kernel_sub(i) によって、kernel_sub　ルーチンに動作が移る。これは、ROUTINE ディレクティブによってデバイス上で動作する対象となっている。デバイス上のグローバル配列 a, b, c を使用して処理される。その後、メインプログラム上で当該 parallel 領域の処理が終了してから update host (c) ディレクティブを用いて、デバイス上の C 配列の内容をホスト側にコピー(更新）する。これにより、ホスト側に計算結果 C 配列の値が戻される。

ソースプログラム : mod.f90、declare_test.f90、kernel_sub.f90、Makefile

mod.f90

module global_data
 implicit none
 integer,parameter :: n= 512
 real(8), dimension(n) :: a, b, c
 !$acc declare create (a,b,c)
 
 interface  ! kernel_subルーチンは別ファイルのため、明示的な interface 指定が必要となる
   subroutine kernel_sub(i)
   !$acc routine vector
     integer :: i
   end subroutine
 end interface
end module global_data

declare_test.f90

program declare_test
use global_data
do i = 1, n
  a(i)=i
  b(i)=1
enddo
! copy a, b in the device memory（実行ディレクティブ）
!$acc update device (a,b)

!$acc parallel loop gang
do i = 1, n
  call kernel_sub(i)
enddo
!$acc end parallel
! copy back c from the device memory（実行ディレクティブ）
!$acc update host (c)

do i = 1, 10
  print *,c(i)
enddo
end program declare_test

kernel_sub.f90

subroutine kernel_sub(i)
!$acc routine vector
use global_data
implicit none
integer :: i
  c(i) = a(i) + b(i)  ! グローバル変数を使う
end subroutine kernel_sub

コンパイル＆実行

$ make
pgf90  -ta=tesla -acc -Minfo=accel -O2 -c mod.f90 -o mod.o
pgf90  -ta=tesla -acc -Minfo=accel -O2 -c declare_test.f90 -o declare_test.o
declare_test:
     10, Generating update device(b(:))
         Generating update device(a(:))
     12, Accelerator kernel generated
         13, !$acc loop gang ! blockidx%x
     12, Generating Tesla code
     19, Generating update host(c(:))
pgf90  -ta=tesla -acc -Minfo=accel -O2 -c kernel_sub.f90 -o kernel_sub.o
kernel_sub:
      1, Generating acc routine vector
         Generating Tesla code
pgf90  -ta=tesla -acc -Minfo=accel -O2 mod.o declare_test.o kernel_sub.o -o a.out
$ ./a.out
    2.000000000000000
    3.000000000000000
    4.000000000000000
    5.000000000000000
    6.000000000000000
    7.000000000000000
    8.000000000000000
    9.000000000000000
    10.00000000000000
    11.00000000000000

declare data ディレクティブ device_resident clause (節)

※以下の仕様の中で common ブロック名を扱えるとの記述がありますが、デバイス側のリンカーの制約で、まだこの用途では common ブロックは利用出来ません。この場合は、common ブロックを使わず、module を使用した形に変更することをお勧めします。

　device_resident clause は、指定した変数用のメモリの確保をアクセラレータ・デバイス・メモリ上のみに行い、ホスト上のメモリ域には確保しないことを指示するものである。この clasue の引数リストに指定された「名前」は、変数名、配列名、スラッシュ記号で囲まれた common ブロック名が許され、サブ配列の指定はできない。ホストは、device_resident clauseの中にある変数にはアクセスできない。device_resident clauseの中で指定されたグローバル変数あるいは common ブロックのアクセラレータ内のデータライフタイムは、プログラムが実行している全ての間となる。

　Fortran において、もし、変数が allocatable 属性を有する変数の場合、当該変数のデバイスメモリ上のメモリ割付と開放のタイミングは、ホストプログラムその変数に対して allocate あるいは deallocate 文の実行が行われた時点となる。もし、変数が Fortran ポインタ属性を有している場合、それはホストによってアクセラレータデバイスメモリ内の割付あるいは開放が行われる。このタイミングは、当該変数がポインタ代入文の左辺側に現れた時点、あるいは、右辺側変数自身が device_resident clause に指定されている時のタイミングとなる。

　Fortran において、device_resident clause への引数がスラッシュ記号で囲まれた common ブロック名の場合もある。このケースでは全ての common ブロック名の宣言は、device_resident clause のもにに合致しなければならない。この場合、 common ブロックは静的にデバイスメモリ内に割り付けられ、ホスト上のメモリには割り付けない。なお、common ブロックは、accelarator routine で利用可能である。

　Fortran Module 宣言部において、device_resident clauseで指定された変数、配列は、accelerator routine で利用可能である。

　C/C++グローバルスコープ内において、device_resident clauseで指定された変数、配列は、accelerator routine で利用可能である。C/C++ extern変数は device_resident clause 内に指定してもよい。その場合は、その変数の実宣言や全ての extern 宣言の後に、device_resident clause を指定する。

※「accelerator routine」とは、アクセラレータ用に routine ディレクティブを指定した C or C++ 関数、Fortran サブプログラムのことを言う。

declare data ディレクティブ link clause (節)

この機能は、現在実装されておりません。

　link clause は accelerator routine 内で参照される、ホスト上の大きなサイズのグローバル静的データのために使用され、デバイス上では動的なライフタイムを有するデータとして利用する。link clause は、名前を有する変数に対するグローバルなリンクだけがアクセラレータのメモリ内に静的に生成されることを指定するものである。ホスト側のデータ構造としてそのグローバル性と静的に割付されたものは、そのまま維持される。デバイス上では、当該データのメモリは、そのグローバル変数が data 構文あるいは、compute 構文、 enter data ディレクティブのいずれかの clause 内に現れた時だけ、割付られる。link clause への引数は、グローバルデータでなければならない。C or C++ では、link clause はグローバル・スコープ上に現れなければならない。また、その引数は extern 変数でなければならない。Fortran では、link cluase は、Moduleの宣言部に現れなければならない。あるいは、その引数は、スラッシュ記号で囲まれた common ブロック名でなければならない。declare link は、当該グローバル変数、あるいは common ブロック変数が、data clause、compute 構文、あるいは accelerator routine　内で明示的にも暗黙的にも使用されるものとして、どこにおいてでも visible なものでなければならない。グローバル変数、あるいは common ブロック変数は、accelerator routine で利用可能である。link clause 内で指定された変数あるいは common ブロックのアクセラレータ上のデータ・ライフタイムは、data cluase を使って変数あるいは common ブロックのエリアをアロケートした時のデータ領域の区間となる。あるいは、 enter data ディレクティブの実行によりデータを割付け、exit data ディレクティブにより割付が開放されるまで、あるいはプログラムの終了ポイントまでの区間となる。

前章へ

次章へ

OpenACCプログラミングのインデックスへ

ページの先頭へ