2

私はOpenCLで作業しています。そして、次の例で作業項目がどのように実行されるかに興味があります。

ワークグループのサイズが 512 で、1 次元の範囲が 10000 です。カーネルは次のとおりです。

__kernel void
doStreaming() {
  unsigned int id = get_global_id(0);

  if (!isExecutable(id))
    return;

 /* do some work */
}

ここでは、次の ID を持つ要素を続行する必要があるかどうかを確認します。

512 サイズの最初のワークグループで実行が開始され、そのうちの 20 が によって拒否されたとしisExecutableます。GPU は最初の 492 要素を待たずに他の 20 要素を実行し続けますか?

関連する障壁やその他の同期技術はありません。

4

2 に答える 2

2

一部のワークアイテムが通常の /* do some work */ から遠く離れている場合、現在のワープ/ウェーブフロント ワークアイテムが他のことを行うのに忙しいため、次のウェーブフロント (amd) または次のワープ (nvidia) から命令を取得することにより、パイプライン占有の利点を利用できます。ただし、これによりメモリ アクセスのシリアル化が発生し、ワークグループのアクセス順序が消去され、パフォーマンスが低下する可能性があります。

ワープ/波面が分岐しないようにします。if ステートメントをループで実行すると、非常に悪いので、別の方法を見つけたほうがよいでしょう。

ワークグループ内のすべてのワークアイテムが同じ分岐を持っていれば問題ありません。

すべての作業項目が数百回の計算でほとんど分岐しない場合は問題ありません。

GPU が持つパワーを利用するために、すべてのワークアイテム (非常に並列なデータ/アルゴリズム) に対して等しい条件を生成するようにしてください。

最も単純な分岐と計算のケースを取り除くために私が知っている最善の方法は、グローバルな yes-no 配列を使用することです。0=はい、1=いいえ : 常に計算し、結果に work-item の yes-no 要素を掛けます。一般に、コアごとに 1 バイト要素のメモリ アクセスを追加することは、コアごとに 1 つの分岐を行うよりもはるかに優れています。実際には、この 1 バイトを追加した後、オブジェクトの長さを 2 の累乗にする方がよい場合があります。

于 2013-07-07T22:57:41.097 に答える
1

はいといいえ。以下の詳細は NVIDIA のドキュメントに基づいていますが、ATI ハードウェアでは異なるとは思えません (実際の数値は異なる可能性があります)。一般に、ワーク グループのスレッドは、ワーク グループ サイズのサブブロックである、いわゆるワープで実行されます。NVIDIA ハードウェアでは、各ワーク グループはそれぞれ 32 スレッドのワープに分割されます。そして、これらの各ワープはロックステップで実行されるため、完全に並列で実行されます (リアルタイムの並列ではない場合があります。つまり、16 のスレッドが並列に実行され、その後すぐに 16 のスレッドが再び実行される可能性がありますが、概念的には完全に並列に実行されています)。 . したがって、これら 32 のスレッドのうちの 1 つだけがその追加コードを実行すると、他のスレッドはそれを待ちます。しかし、他のすべてのワープのスレッドは、これらすべてを気にしません。

そうです、他のスレッドを不必要に待機するスレッドが存在する可能性がありますが、それはワーク グループ全体のサイズ (NVIDIA ハードウェアでは 32) よりも小さい規模で発生します。これが、可能であればワープ内の分岐逸脱を回避する必要がある理由であり、単一のワープ内でのみ動作することが保証されているコードが、共有メモリ アクセスなどの同期を必要としない理由でもあります (アルゴリズムの一般的な最適化)。

于 2013-07-07T23:04:20.737 に答える