15

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_elementsgentype要素の非同期コピーを実行します。非同期コピーはによって実行されます。したがって、ワークグループ内のすべてのワークアイテムとこの組み込み関数は、同じ引数値でカーネルを実行しているワークグループ内のすべてのワークアイテムが遭遇する必要があります。そうでない場合、結果は未定義です。」

しかし、それは私の質問を明確にしません...

次の仮定が正しいかどうか知りたいのですが。

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

ボーナスの質問：

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

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

よろしく、ステファン

4

1 に答える 1

12

この関数が存在する主な理由の1つは、開発者がハードウェアについて想定することなく、ドライバー/カーネルコンパイラーがメモリを効率的にコピーできるようにすることです。

コピーする必要のあるメモリをシングルスレッドのコピーであるかのように記述し、async_work_group_copyが並列ハードウェアを使用してそれを実行します。

あなたの特定の質問のために：

  1. グループ内の一部の作業項目でのみasync_work_group_copyが使用されているのを見たことがありません。必要だったからだといつも思っていました。wait_group_eventsのブロックの性質により、すべての作業項目がコピーの一部になると思います。

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

  3. ローカルIDを差し引いて正しいアドレスを取得することもできますが、アドレスをgroupIdに基づいて計算すると、この問題も解決されることがわかりました。（get_group_id）

  4. はい。最後のパラメータは要素の数であり、バイト単位のサイズではありません。

a。いいえ。イベントベースでは、ほとんどすぐに作業項目が障壁にぶつかり、データが必ずしもコピーされるとは限りません。一部のopenclハードウェアは、実際のコピー操作を行うために計算ユニットをまったく使用しない可能性があるため、これは理にかなっています。

b。cpu openclの実装は、ローカルメモリを使用するときにL1キャッシュの使用を保証する可能性があると思います。これがより良いパフォーマンスを発揮するかどうかを確実に知る唯一の方法は、さまざまな設定でアプリケーションをベンチマークすることです。

于 2013-03-21T13:10:49.100 に答える