Half-Wrapってどう作られる？

要はどれとどれのスレッドが1つのHalf-Wrapになるか誰か知りませんかね。
まぁCUDAの話なんですけど。

coalesced accessになるように最適化しようと思ってみたものの、どういう要件でHalf-Wrapにまとめられるかが分からず・・・
やっぱ blockDim.x*blockDim.y*threadIdx.z + blockDim.x*threadIdx.y + threadIdx.x で決まるのだろうか。また仕様書を読み直す作業が・・・
英語あんま得意じゃないんだよなぁ。