openacc - async - 多くの非同期キューを使用するとデータ転送速度が向上する

Question

私の質問は、ホストとデバイス間の異なる非同期キューで複数のアレイを転送する影響についてです。

double *a, *b, *c, *d; 

そして、それぞれがNのサイズで割り当てられています

は、我々は4つの配列を持っていると仮定します。我々は多くの非同期条項にを使用し、異なるキューにそれらを転送することができ

#pragma acc enter data copyin(a[0:N], b[0:N], c[0:N], d[0:N]) async 
#pragma acc wait 

または、：：

a = (double*) malloc(N * sizeof(double)); 
b = (double*) malloc(N * sizeof(double)); 
c = (double*) malloc(N * sizeof(double)); 
d = (double*) malloc(N * sizeof(double)); 

今、私たちは、1節に、次のようなデバイスとホストとの間でそれらを転送することができます

#pragma acc enter data copyin(a[0:N]) async(1) 
#pragma acc enter data copyin(b[0:N]) async(2) 
#pragma acc enter data copyin(c[0:N]) async(3) 
#pragma acc enter data copyin(d[0:N]) async(4) 
#pragma acc wait 

上記の両方のアプローチの結果は同じです。しかし、パフォーマンス面では、2番目の方が良い場合もあります。

私はいくつかの測定を行いましたが、ホストの更新と更新では、moreを使用しているキューの方が性能面ではoneより優れているようです。

のは、最初のアプローチone、及び第二のアプローチmore、およびアプローチmore_nonumber次（非同期句には数に気付かない）を呼び出してみましょう：次に

#pragma acc enter data copyin(a[0:N]) async 
#pragma acc enter data copyin(b[0:N]) async 
#pragma acc enter data copyin(c[0:N]) async 
#pragma acc enter data copyin(d[0:N]) async 
#pragma acc wait 

、ここでは（最初の100を除く万回の繰り返しのための測定であります間に9,800回の反復の平均を導く100の最後のもの）：

one

COPYIN：64.273us

更新デバイス：60.928us

更新自己：69.502us

CopyOut：70.929us

---

more

COPYIN：65.944us

更新装置：62.271us

更新自己：60.592us

CopyOut：59.565us

---

more_nonumber

COPYIN：66。018us

更新デバイス：62.735us

更新自己：70.862us

CopyOut：72.317us

---

平均9800の実行！

one、または更新自己（69.502/60.592）の14％と比較しmore方法を使用する場合、19％の高速化は、（70.929/59.565）copyout観察されます。

私の質問： これらの数字は正当なものですか？これらの数字に頼ることはできますか？

お客様の便宜のため、githubに私のコードを記載しました。あなたはそれを見ることができます。

Answer 1

非同期は、データの移動と計算をデバイスでインターリーブするときに最も便利です。したがって、この演習はちょっと関係がありませんが、私は何が起こっているのかを説明するために最善を尽くします。私はこれが現在PGI（v16.10）が "非同期"を実装しており、他のOpenACC実装がどのように "非同期"を実装するかということではないことに注意してください。

デフォルトでは、PGIはダブルバッファリングシステムを使用してデータ転送を実行します。 DMA転送は物理メモリに存在しなければならないため、バッファは固定されています。ランタイムは仮想メモリを固定バッファにコピーし、バッファの非同期転送を開始し、次に第2バッファの仮想から固定されたコピーを開始します。バッファはいっぱいになり、次にフルアレイがコピーされるまで順番に転送されます。各非同期キューには独自のバッファがあります。

非同期節を使用せずにコードを実行すると、4つの変数を1つのプラグマに入れることが、それぞれの変数を持つことが非常に高速になることがわかります。その理由は、4つのプラグマを使用すると、ホストは次の変数に移動する前に最後のバッファが送信されるまで待機します。それらがすべて同じプラグマ内にある場合、配列の最後のバッファが開始されると転送が開始され、ランタイムは次の配列のデータでもう一方のバッファを埋めることができます。

単一のプラグマに「async」を追加してから待機すると、「async」または「wait」をまったく使用しない場合とパフォーマンスの違いは見られません。言い換えれば、これらは同じです：あなたはCPUが同じプラグマでそれらすべてを入れてしまうよう

#pragma acc update device(a,b,c,d) 

#pragma acc update device(a,b,c,d) async 
#pragma acc wait 

あなたは4つの個々のプラグマに「非同期」を追加し、あなたが同じパフォーマンスについて得られますdoesnの次の配列のバッファリングを開始するのを待ちます。

各アレイが独自の非同期キューにあるとき、デバイスからのコピーバック（自己、コピーアウトの更新）が高速である理由を説明できません。それは大したことではないと私には思われます。

このテストでは、最適なのはダブルバッファを使用せず、代わりに物理メモリにアレイ全体を固定することです。これにより、仮想メモリから固定メモリにコピーするコストを節約できます。物理メモリは限られており、プログラムのメモリが適切であるとは保証できないため、コンパイラはデフォルトでこれを行うことはできません。ほとんどのコードでは、ピンメモリを使用してデバイス上で計算するのに費やされていても、あまり効果がありません。最後に、固定メモリの割り当てを解除するためのパフォーマンスオーバーヘッドがあります（固定メモリへの転送を保証するためにデバイスを同期させる必要があります）。したがって、配列が単一のデータ領域にあるが、多くの更新がある場合は、最も有益です。

ここでは固定メモリを使用した回です：包括的な対応のための

% pgcc -fast -acc -ta=tesla:cc35,pinned transfer.c -o tpinned.out 
% numactl -C 2 ./tpinned.out 
one 
CopyIn: 50.161us 
Update device: 49.679us 
Update self: 47.595us 
CopyOut: 47.631us 
--------- 
more 
CopyIn: 52.448us 
Update device: 52.135us 
Update self: 49.904us 
CopyOut: 47.926us 
--------- 
more_nonumber 
CopyIn: 52.172us 
Update device: 51.712us 
Update self: 49.363us 
CopyOut: 49.430us 
---------