OpenCLでasync_work_group_copyを使用するには？

Question

OpenCLでasync_work_group_copy（）を正しく使用する方法を理解したいと思います。のは、簡単な例で見てみましょう：。

__kernel void test(__global float *x) { 
    __local xcopy[GROUP_SIZE]; 

    int globalid = get_global_id(0); 
    int localid = get_local_id(0); 
    event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0); 
    wait_group_events(1, &e); 
} 

参照がhttp://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html srcからdstへNUM_ELEMENTS gentype要素の非同期コピーを実行します」と言うの非同期コピーがワークグループ内のすべての作業項目によって実行され、したがって、この組み込み関数は、同じ引数値でカーネルを実行しているワークグループ内のすべての作業項目によって検出されなければなりません。そうでなければ、結果は未定義です。

しかし、それは私の質問を明確にしない...

次の仮定が正しければ、私は、知っているしたいと思います：

1. （async_work_group_copyの呼び出し）は、すべての作業によって実行されなければなりませんグループ内の項目。
2. 呼び出しは、送信元アドレスがすべての作業項目で同一であり、コピーするメモリ領域の最初の要素を指している必要があります。
3. ソースアドレスは、ワークグループ内の最初の作業項目のグローバル作業項目IDに基づいているため、だから私はすべての作業項目のアドレスを同一にするためにローカルIDを引く必要があります...
4. 3番目のパラメータは本当に要素の数です（バイト単位ではありません）？

ボーナス質問：

。 wait_group_events（）の代わりにバリア（CLK_LOCAL_MEM_FENCE）を使用して戻り値を無視することはできますか？もしそうなら、それはおそらくもっと速いでしょうか？

b。ローカルコピーはCPU上での処理にも意味がありますか、それともキャッシュを共有しているのでオーバーヘッドですか？

よろしく、 ステファン

Answer 1

既存この機能の主な理由の一つは、ドライバ/カーネルのコンパイラは、効率的なハードウェアについての仮定を作るために持っ開発せずにメモリをコピーできるようにすることです。

コピーする必要があるメモリをシングルスレッドコピーのように記述し、async_work_group_copyを使用すると、並列ハードウェアを使用して処理されます。あなたの具体的な質問について

：

1. 私は、グループ内の作業項目の一部だけが使用するasync_work_group_copyを見たことがありません。私はいつもそう思っていました。私は、wait_group_eventsの性質を阻止することによって、すべての作業項目がコピーの一部になるようにしていると思います。

2. はい。ソース（および宛先）アドレスは、すべての作業項目で同じである必要があります。

   ローカルIDを減算して正しいアドレスを取得できますが、groupIdのアドレスを基にすればこの問題も解決されます。 （get_group_id）

3. はい。最後のparamはバイト数ではなく要素数です。

a。いいえ。イベントベースでは、障壁が作業項目によってほぼ直ちに打たれ、必ずしもデータがコピーされるとは限りません。オープンループのハードウェアの中には、実際のコピー操作を行うために計算ユニットをまったく使用しないものもあるかもしれないので、これは理にかなっています。

b。私は、ローカルメモリを使用すると、CPUキャッシュのopenclの実装でL1キャッシュの使用が保証されると思います。これが優れているかどうかを確認する唯一の方法は、さまざまな設定でアプリケーションをベンチマークすることです。