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

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

9章　その他のディレクティブ

Declare ディレクティブ

　declare ディレクティブは、 Fortran サブルーチン、関数、モジュール（MODULE)の宣言部、または、C/C++の変数宣言の直後に指定する。これは、関数、サブルーチン、プログラムの暗黙のデータ領域とされる場所で、変数あるいは配列がデバイスメモリ上にアロケートされることを指示するために使用する。そして、暗黙のデータ領域の開始するポイントにおいて当該データがホストからデバイスへのデータコピーを行うかどうか、あるいは、暗黙のデータ領域の終了ポイントでデバイスからホストへデータをコピーするかどうか等の挙動も指定することができる。このディレクティブは、変数あるいは配列の visible device copy (コンパイラがコンパイル段階でコピー挙動が明示的に分かること）の生成を行う。

【Syntax】
C and C++ の declare ディレクティブ
　 #pragma acc declare clause-list new-line 
 
Fortran の declare ディレクティブ
　 !$acc declare clause-list 

Clause（節）は、次のものが指定できる
 copy( var-list )
 copyin( var-list )
 copyout( var-list )
 create( var-list )
 present( var-list )
 present_or_copy( var-list )
 present_or_copyin( var-list )
 present_or_copyout( var-list )
 present_or_create( var-list )
 deviceptr( var-list )
 device_resident( var-list )
 link( var-list )　[OpenACC 2.0]

該当する領域は、当ディレクティブが現れている関数、サブルーチン、プログラムに関連した暗黙のデータ領域に対してとなる。もし、このディレクティブが Fortran Module サブプログラムの宣言部、あるいは、C/C++ のグローバルなスコープエリア内に指定されると、その対応する領域はプログラム全体が暗黙のデータ領域とみなされる。それ以外、当該 clause (節) は、これら clause を使ってプロシジャー本体に対して適用する「明示的な」data 構文と全く同じように振る舞う。

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 だけが許される。（OpenACC 2.0)
C/C++のグローバル・スコープ下では、create, copyin, deviceptr, device_resident とlink の clause だけが許される。（OpenACC 2.0)
C/C++の extern 変数は、create, copyin, deviceptr, device_resident とlink の clause 内で指定できる。（OpenACC 2.0)

Combined(複合)ディレクティブ

　複合 Parallel loop と kernels loop ディレクティブは、これに続くネストループに対して parallel あるいは kernels 構文の内部で loop ディレクティブを指定している状態と同じ機能を提供する。すなわち、明示的に loop ディレクティブを含んだ parallel あるいは kernels ディレクティブを指定していることと同じ意味を成す。この Parallel loop ディレクティブで利用できる clause は、parallel あるいは loop ディレクティブで許されているものと同じとなる。同様に kernels loop ディレクティブで利用できる clasue は、kernels あるいは loop ディレクティブで許されているものと同じとなる。このディレクティブは、これに続く直後のループに対して機能する。このディレクティブの使用例は、8章の姫野ベンチマークのプログラムにあるので、参照して欲しい。

【Syntax】
C and C++ の parallel loop ディレクティブ
　 #pragma acc parallel loop [clause-list] new-line 
        for loop
 
Fortran の parallel loop ディレクティブ
　 !$acc parallel loop [clause-list] 
　      do loop
　 !$acc end parallel loop 

C and C++ の kernelsl loop ディレクティブ
　 #pragma acc kernels loop [clause-list] new-line 
        for loop
 
Fortran の kernels loop ディレクティブ
　 !$acc kernels loop [clause-list] 
　      do loop
　 !$acc end kernels loop

制限事項

parallel、kernels、loop 構文の制限事項が適用される。

update 実行ディレクティブ

　update ディレクティブは、アクセラレータのライフタイムの間に、デバイス・メモリ内の変数、配列に対応するローカル（ホスト）側変数、配列の一部あるいは全部の値を更新する機能、あるいは、ローカル（ホスト）側のメモリ内の変数、配列に対応するデバイス上の変数、配列の一部あるいは全部の値を更新するための機能を提供する。これは、「実行文」として機能するため、ディレクティブを挿入した箇所で、その機能が実行される。update ディレクティブを使用した例を第12章の「OpenACC Fortran と CUDA library (cufft)」に示したので参照のこと。

【Syntax】
C and C++ の update ディレクティブ

　 #pragma acc update [clause-list] new-line 
 
Fortranの update ディレクティブ

　 !$acc update [clause-list] 

Clause（節）は、次のものが指定できる
　  async [( int-expr )] 
    wait [( int-expr-list )] 
    device_type( device-type-list ) 
    if( condition ) 
    self( var-list ) 
    host( var-list ) 
    device( var-list )

　update ディレクティブの clause の var-list 引数には、コンマで区切って変数、配列、サブ配列を指定する。同じ配列の中のサブ配列を複数指定できる。すなわち、同じディレクティブの上でこれらは同じ clause 内、もしくは異なる clause の var-list の中に現れてもよい。update の clause の機能として、update self とするとアクセラレータ・デバイスからローカル（ホスト）・メモリへのコピーを行い、一方、update device とするとローカル・メモリからデバイス・メモリへのコピーを行う。その更新は、ディレクティブ上に指定した順番で行われる。このディレクティブを使用するには、少なくとも一つの self、host あるいは device clause を指定する必要がある。

self clause

　self clause は var-list 内に指定された変数、配列、サブ配列が、アクセラレータ上のデバイスメモリからローカルメモリへコピーすることを指示するものである（non-shared メモリ型のアクセラレータ構成の場合）。もし、アクセラレータ上で当該スレッドがそれぞれ同じメモリを共有している場合、このアクションは実行されない。なお、「ローカルメモリ」の定義に関しては、7章を参照いただきたい。

host clause

　host は、self と同義語である。これは、OpenACC 1.0 において使用されていた clause である。OpenACC 1.0 では、デバイスメモリからホストメモリへデータコピーを行うための指定方法である。

device clause

　device clause は var-list 内に指定された変数、配列、サブ配列が、ローカルメモリからアクセラレータ上のデバイスメモリへコピーすることを指示するものである（non-shared メモリ型のアクセラレータの場合）。もし、アクセラレータ上で当該スレッドとそれぞれ同じメモリを共有している場合、このアクションは実行されない。

if clause

　if clause はオプションである。この if clause が指定されていない場合、コンパイラは、アップデートを実行するコード生成する。if clause を指定されている場合は、条件付きでアップデートを行うコードを生成する。この条件の評価法は、C/C++ の場合はノンゼロで Fortran の場合は .true. の時のみ、更新が実行される。

async と wait clause

　6 章を参照されたし。あるいは、OpenACC 2.0 の仕様書を確認してください。

制限事項

update ディレクティブは実行文である。これは、C/C++ 言語の場合は if, while, do, switch あるいは labelの後、Fortran 言語の場合は論理 if 文の次の場所では指定してはいけない
update ディレクティブの var-list に指定されている変数、配列はそのデバイス上のコピーを有していなければならない
async と wait clauses だけは、a device_type clause の後になければならない
if clause がある場合、Fortran での条件はスカラ論理値で評価する。C/C++ の条件はスカラ整数値で評価する。
非連続のサブ配列が指定された場合の挙動は、コンパイラの実装で異なるかもしれない。非連続の領域に関して、各連続したサブ領域の各々を 1 回づつ転送を行うことによって更新するのか、あるいは、非連続データをパックして、1 回の転送を行い、その後、アンパックするような手法をとるかは、コンパイラによって異なる。
C/C++ では、struct あるいは class のメンバーがそのメンバーのサブ配列を含んで指定される場合もある。struct あるいは class のサブ配列のメンバーを指定してはならない
C/C++ では、サブ配列の notation が、struct のメンバーに対して使用されている場合、サブ配列の notation はその struct のメンバーの親に対して使用してはならない
Fortran では、派生型変数のメンバーがそのメンバーのサブ配列を含んで指定される場合もある。派生型のサブ配列のメンバーを指定してはならない
Fortran ではサブ配列の notation が、派生型のメンバーに対して使用されている場合、配列あるいはサブ配列の notation はその派生型のメンバーの親に対して使用してはならない