2

CUDAプログラミングガイド（v4.1）は、セクション5.4.2の述語命令についてこれを説明しています。

コンパイラは、分岐条件によって制御される命令の数が特定のしきい値以下である場合にのみ、分岐命令を述語命令に置き換えます。条件が多くの発散ワープを生成する可能性があるとコンパイラが判断した場合、このしきい値は7です。 4です。

  1. 条件はどのようにして多くの発散ワープを生成することができますか？特定の条件では、ワープを2つに分割することしかできません。ここで多くの意味は何ですか？
  2. 上記が理にかなっているとしても、コンパイラはワープの実行時の発散動作をどのように知ることができますか？
4

1 に答える 1

2

ワープが「分割」されることはありません。それらは、条件付きで発散するコードパスを処理するために「条件付き実行」（つまり、参加していないスレッドをマスクした状態での実行）を必要とするか、または必要としません。

条件が複数の発散ワープを生成する方法については、次の不自然な例を検討してください。

if (threadIdx.x < 128) {
   // Only first four warps process here
   int modthirtytwo = threadIdx.x % 32;

   if (modthirtytwo == 0) {
      // Action A only first thread in the warp
   } else {
      // Action B for the other threads in the warp
   }
}

ここで、コードは複数の発散ワープを生成する可能性があり、コンパイラーはコンパイル時に動作をモデル化できる必要があります。カーネルのコンパイラーに起動境界が指定されていれば、さらに良いでしょう。このケースを、ワープを1つだけ使用する共有メモリの削減と比較してください。

if (threadIdx.x < 32) {
   if (threadIdx.x < 16)  shm[threadIdx.x] += shm[threadIdx.x+16];
   if (threadIdx.x < 8)   shm[threadIdx.x] += shm[threadIdx.x+8];
   if (threadIdx.x < 4)   shm[threadIdx.x] += shm[threadIdx.x+4];
   if (threadIdx.x < 2)   shm[threadIdx.x] += shm[threadIdx.x+2];
   if (threadIdx.x == 0)  shm[0] += shm[1];
}

ここでは、発散はブロックごとに1つのワープに制限されています。そのテキストのすべては、2つの場合のコンパイラの動作が異なる可能性があるということです。

「新しい」コンパイラー（OpenCLで数年間使用されています）には、ブランチがより経済的になる前に使用する必要のある述語命令の数についてのヒューリスティックがあるようです。また、命令パイプラインの多くのブランチはパフォーマンスに良くないようです。そのため、コンパイラがそのコードがより高い「ブランチ密度」を生成することを理解できる場合、ブランチではなく、より多くの述語命令を優先します。

于 2012-04-12T09:01:48.210 に答える