CUDAコンパイラはどのようにワープの発散動作を知っていますか？

Question

CUDAプログラミングガイド（V4.1）を説明について断定指示秒5.4.2でこのコンパイラは、命令の数によって制御する場合にのみ前提 命令で分岐命令を置き換え

分岐条件が一定のしきい値に小さいか等しい： コンパイラは条件が多く 発散縦糸を生産する可能性があると判断した場合、このしきい値は7で、それ以外の場合は、4

です

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

Answer 1

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

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に数年間使用されています）は、分岐がより経済的になる前にいくつの述語命令を使うべきかについての経験を持っているようです。そして、命令パイプラインの多くの分岐がパフォーマンスに悪いように見えるので、コンパイラがコードがより高い "分岐密度"を生成するようになると、分岐の代わりにより多くの述語命令が優先されます。