1

一部のデータを共有メモリに書き込むワープがあります-上書きせずに、共有メモリから読み取った直後に。私のブロックには他のワープがあるかもしれませんが、それらはその共有メモリのどの部分にも触れず、関心のあるワープが読み取る場所に書き込みません。

ここで、ワープがロックステップで実行されているにもかかわらず、共有メモリの書き込みに続く共有メモリの読み取りが、ワープによって以前に書き込まれたと思われるそれぞれの値を返すことが保証されていないことを思い出します。(これは、理論的には、命令の並べ替えが原因である可能性があります。または、@RobertCrovella が指摘するように、共有メモリ アクセスを最適化するコンパイラが原因である可能性があります)

したがって、明示的な同期に頼る必要があります。明らかに、ブロックレベルの__syncthreads()作業です。これは何をするかです：

__syncthreads()同じブロックのスレッド間の通信を調整するために使用されます。ブロック内の一部のスレッドが共有メモリまたはグローバル メモリ内の同じアドレスにアクセスする場合、これらのメモリ アクセスの一部について、書き込み後読み取り、読み取り後書き込み、または書き込み後書き込みの危険が生じる可能性があります。これらのデータ ハザードは、これらのアクセスの間にスレッドを同期することで回避できます。

それは私のニーズにはあまりにも強力です:

  • 共有メモリだけでなく、グローバルメモリにも適用されます。
  • ワープ間の同期を実行します。イントラワープのみが必要です。
  • R-after-W、W-after-R、W-after-W のすべてのタイプの危険を防ぎます。R-after-Wだけが必要です。
  • 複数のスレッドが共有メモリ内の同じ場所への書き込みを実行する場合にも機能します。私の場合、すべての共有メモリへの書き込みはばらばらです。

一方、のようなもので__threadfence_block() は十分ではないようです。これら 2 つのレベルの強さの「中間」にあるものはありますか?

ノート：

  • 関連する質問: warp 内でのCUDA__syncthreads()の使用。
  • 代わりにシャッフルを使用することを提案する場合は、はい、それが可能な場合もありますが、データへの配列アクセスが必要な場合、つまり、共有データのどの要素を読み取るかを動的に決定する場合はそうではありません。それはおそらくローカルメモリに流出するでしょう、それは私には恐ろしいようです.
  • 私はおそらくvolatile私に役立つかもしれないと思っていましたが、それを使用することで私が望むことができるかどうかはわかりません.
  • コンピューターの機能が少なくとも XX.YY であると想定する回答があれば、それで十分です。
4

1 に答える 1

1

@RobertCrovellaを正しく理解していれば、このコードの断片は危険から守られているはずです:

/* ... */
volatile MyType* ptr = get_some_shared_mem();
ptr[lane::index()] = foo();
auto other_lane_index = bar(); // returns a value within 0..31
auto other_lane_value = ptr[other_lane_index];
/* ... */

の使用のためvolatile。bar()（そして、混乱しないと仮定することは、それ自体の危険をもたらします。）

于 2017-04-20T13:32:44.857 に答える