OpenACC ディレクティブによるプログラミング
6章 Accelerator Compute 構文
kernels 構文
kernels 構文についての概説は、前章で述べた。kernels 構文は、アクセラレータ上の（一連の）「カーネル」として、コンパイラがコード生成を行うプログラムの領域を指定するためのディレクティブである。
kernels ディレクティブを使って、アクセラレータ上で並列実行させたい「tightly nested loop」の構造化ブロックに対してその領域を指定する。以下の図の例では、対象とするブロックは三重のネストループとその後にシングルループと言う形で構成されているが、この全体を一つの kernels 構文の領域として指定している。こうした場合、コンパイラは上の三重ループを一つのカーネル (kernel1) として生成し、次のシングルループをもう一つ独立のカーネル (kernel2) として生成する。実行時 kernels 構文の場合は、kernel1 の実行が終了してから kernel2 の実行を始める。必ず、カーネルのスレッド実行が全て終了時点で、次のカーネルを実行すると言う同期を取る。これは kernels 構文の特徴と言える。次に説明する parallel 構文の場合は、その中に複数のループが存在する場合、このような全スレッドの終了を待って次のループの開始を行うと言うことは一切しない。この違いを良く理解しておく必要がある。
kernels ディレクティブのシンタックスと clauses（節）
kernels 構文のシンタックスと利用可能な clause（節）を以下に示した。
【Syntax】 C and C++の場合 #pragma acc kernels [clause-list] new-line { 構造化ブロック（ループ）} Fortranの場合 !$acc kernels [clause-list] 構造化ブロック（ループ） !$acc end kernels Clause（節）は、次のものを指す async [( int-expr )] wait [( int-expr-list )] device_type( device-type-list ) : OpenACC 2.0 if( condition ) 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 ) default( none ) : OpenACC 2.0
async 節が指定されていない場合は、kernels 領域の最後の時点で暗黙のバリア同期を行う。また、全ての「カーネル」が kernels 領域の終点で終了するまでホスト（ローカル）側のスレッドの実行は行わない。
default(none) 節が指定されていない場合、コンパイラは、以下のものを除いて Compute 構文内で使用されている変数に対して暗黙にデータ（転送）属性を決定する。
- Compute 構文の data clause で明示的にデータ属性を指定している変数
- Compute 構文の前に、data 構文で明示的にデータ属性を指定している変数
上記の状態の「配列あるいは集合データ型」は、コンパイラによる暗黙に指定される属性としては present_or_copy が設定される。また、同様にスカラ変数に関しては copy 属性を設定する。なお、present_or_* clause の詳細は、第7章 data 構文の clauses の項を参照のこと。
制限事項
- プログラムは、kernels 構文ループ内の途中でループ外へジャンプすることを含めて、exit しないこと。また、その逆も許されない。すなわち、ループは「構造化ブロック」であること。
- kernels 構文の clause の指定順序によって、その評価の優先度が左右されることはない（指定順序は関係ない）
- async と wait cluase だけは、device_type clause の後に置く必要がある（device_type を指定した場合）
- if clause は、Fortran ではその条件はスカラ論理値として評価する必要がある。C/C++ ではスカラ整数値として評価する必要がある
parallel 構文
parallel 構文についての概説は、前章で述べたparallel 構文は、アクセラレータ上で並列実行を開始させるためのディレクティブである。
プログラムが parallel 構文に到達した時点で、アクセラレータ上で当該 parallel 領域を実行するために、一つ以上の gang が生成される。その際、gang の数、gang 当たりの worker の数、そして worker 当たりの vector lane の数は、その parallel 領域の実行の間、一定に維持される（増減されない）。各 gang はプログラムの構造化ブロック内で gang-redundant モードでコードを実行し始める。このことは、loop ディレクティブを指定したループ、あるいは gang-level work-sharing が指定されたループ以外の parallel 領域内のコード部分は、全ての gang によって冗長に実行されることを意味している。真の「並列実行」を行う部分は、loop ディレクティブを指定して work-sharing を指示したループとなることに注意が必要である。
parallel ディレクティブのシンタックスと clauses（節）
parallel 構文のシンタックスと利用可能な clause（節）を以下に示した。
【Syntax】 C and C++の場合 #pragma acc parallel [clause-list] new-line { 構造化ブロック（ループ）} Fortranの場合 !$acc parallel [clause-list] 構造化ブロック（ループ） !$acc end parallel Clause（節）は、次のものを指す async [( int-expr )] wait [( int-expr-list )] num_gangs( int-expr ) num_workers( int-expr ) vector_length( int-expr ) device_type( device-type-list ) : OpenACC 2.0 if( condition ) reduction( operator : var-list ) 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 ) private( var-list ) default( none or present ) : OpenACC 2.0
async 節が指定されていない場合は、parallel 領域の最後の時点で暗黙のバリア同期を行う。また、全ての「カーネル」が parallel 領域の終点で終了するまでホスト（ローカル）側のスレッドの実行は行わない。
default(none) 節が指定されていない場合、コンパイラは、以下のものを除いて Compute 構文内で使用されている変数に対して暗黙にデータ（転送）属性を決定する。
- Compute 構文の data clause で明示的にデータ属性を指定している変数
- Compute 構文の前に、data 構文で明示的にデータ属性を指定している変数
上記の状態の「配列あるいは集合データ型」は、コンパイラによる暗黙に指定される属性としては present_or_copy が設定される。また、同様にスカラ変数に関しては firstprivate 属性を設定する。なお、present_or_* clause の詳細は、第7章 data 構文の clauses の項を参照のこと。
制限事項
- プログラムは、paralle 構文ループ内の途中でループ外へジャンプすることを含めて、exit しないこと。また、その逆も許されない。すなわち、ループは「構造化ブロック」であること。
- kernels 構文の clause の指定順序によって、その評価の優先度が左右されることはない（指定順序は関係ない）
- async と wait cluase だけは、device_type clause の後に置く必要がある（device_type を指定した場合）
- if clause は、Fortran ではその条件はスカラ論理値として評価する必要がある。C/C++ ではスカラ整数値として評価する必要がある
Accelerator Compute 構文の clauses（節）の説明
async clause と wait clause
async cluase は parallel 構文、kernels 構文、enter data、exit data、update あるいは wait ディレクティブと共に指定可能な clause である。これら全てのケースで、この clasue はオプションである。 async が指定されていない場合、ローカル（ホスト）スレッドは、当該 compute 構文あるいは data 処理が終了するまで、次のコードを実行することを待つ。wait ディレクティブの場合もこれと同じ動作を行う。wait ディレクティブの場合は、対応する非同期実行キューの全ての処理が完了するまで待つ。一方、async clause があるとき、ローカルスレッドは非同期に parallel 構文、kernels 構文、data 処理の次に置かれているコード部分の処理を継続する。さらに詳細は、OpenACC 仕様書を参考にすること。
if clause
if clause は、必須ではなくオプションである。parallel 構文、kernels 構文上で指定できる。if が現れるとコンパイラは、当該構文対象部に関して二つのコピーを生成する。一つは、アクセラレータ上で実行可能なコード、他方はローカルスレッド上で実行可能なコードである。C/C++ の場合は、その条件判定においてノンゼロである場合、Fortran では .true. である場合、アクセラレータ上のコピーが実行される。その逆の条件判定の場合は、ローカルスレッド上のコードが実行される。
num_gangs clause
num_gangs clause は、parallel 構文のみに許される。整数（式）で与えられる数字は、parallel 領域を実行する並列 gang の数を定義する。もし、この数が与えられていない時は、実装依存となりそのデフォルトが使用される。このデフォルト値は、構文内のコードに依存する。
num_workers clause
num_workers clause は、parallel 構文のみに許される。整数（式）で与えられる数字は、worker-single モードから worker-partitioned モードに移行した gang がアクティブになった後、各 gang 内の worker の数を定義する。もし、この数が与えられていない時は、実装依存となりそのデフォルトが使用される。このデフォルト値は、1 もしくは、各 parallel 構文によって異なるかもしれない。
vector_length clause
vector_length clause は、parallel 構文のみに許される。整数（式）で与えられる数字は、 vector-single モードから vector-partitioned モードに移行した worker がアクティブになった後の vector lane の数を定義する。この clause は、vector あるいは SIMD 処理のベクトル長を決定するものである。もし、この数が与えられていない時は、実装依存となる。この ベクトル長は、loop ディレクティブ上の vector clause を指定したループに対して適用される他、コンパイラによる自動ベクトル化されるループに対しも適用される。
（PGI 15.1 追加）以前のバージョンでは、parallel directive の vector_length clause 内の表現式や kernel region 内のloop directive の vector clause 内の表現式は、コンパイル時に分かるように定数を使った表現式であることの制約があった。PGIコンパイラは、この clause 内のランタイム表現式を使えるように変更した。もし、その値がテーゲットとなるデバイスの上限を超えた場合、プログラムはランタイムでエラーとなる。現在のNVIDIA CUDAデバイスの上限値は、vector lengthとworkers の数の積として 1024 が上限となる。AMD Radeon デバイスの上限値は、vector length と workers の数の積として256が上限となる。
private clause
private clause は、parallel 構文のみに許される。コンパイラは、clause のリスト上に記載された変数のコピーが各 parallel gang に対して生成する。各 gang のプライベートな変数としてコード生成される。
first private clause
first private clause は、parallel 構文のみに許される。コンパイラは、clause のリスト上に記載された変数のコピーが各 parallel gang に対して生成する。各 gang のプライベートな変数としてコード生成される。そして、そのコピーは、parallel 構文に到達した際にホスト上の変数が有している値に初期化される。
reduction clause
reduction clause は、parallel 構文のみに許される。これは、一つ以上のスカラ変数のリダクション処理を指定するために使用する。各変数は、各 parallel gang に対してプライベートな変数を作成し、その処理のための初期化を行う。その後、parallel 領域の終了時点で各 gang のその値はリダクション処理により全て集合化される。その集合化された値は、オリジナル変数の値としてストアされる。リダクション結果は、parallel 領域の後で有効となり、使用可能となる。以下に、リダクション可能な処理を列挙した。なお、C/C++ においては、int, float, double, complex データ型の演算がサポートされる。また、Fortran では、integer, real, double precision, complex のデータ型がサポートされている。reduction 節の指定の例は、例えば、sum リダクションであれば reduction(+:val) と言う形態になる。
C and C++ Fortran
-----------------------------------------------------------
operator initialization operator initialization
value value
-----------------------------------------------------------
+ 0 + 0
* 1 * 1
max least max least
min largest min largest
& ～0 iand all bits on
| 0 ior 0
^ 0 ieor 0
&& 1 . and. .true.
|| 0 .or. .false.
.eqv. .true.
.neqv. .false
----------------------------------------------------------
device_type clause (OpenACC 2.0)
device_type clause は、オプションである。この clause を使うことにより複数の異なるアクセラレータに対して特有な clause を定義することが出来る。device_type clause への引数の指定方法は、アクセラレータ名をコンマで区切り指定する。あるいは、任意のアクセラレータという意味で、アスタリスクで指定する。
#pragma acc parallel device_type(nvidia) num_gangs(200) device_type(radeon) num_gangs(400)
default clause (OpenACC 2.0)
default(none) clause は、オプションである。これは、コンパイラに対してデータの属性を暗黙に決定しないように指示するものである。この対象は、computer 構文上で明示的に data clause で指定したものだけでなく、computer 構文上で、自動的に解釈されたデータ、さらに declare ディレクティブで指定された変数も含む。また、default(present) clause 指定した場合は、compute 構文で使用している配列や集合データ内の変数に対して、あたかもこれらが present clause が暗黙に指定されたものとして扱われる。
data clause（コピーの方法を指示する clauses）
copy、cpoyin 等のデータのコピー、アロケートに関する属性は、次章で述べる deta 構文の data clause の意味と同じであるため、次章を参考されたし。
