我知道对于具有31个线程的1D线程块,它将被填充到32个线程以进行warp执行.具有31*31线程的2D块怎么样？warp scheduler会为每个维度填充1个额外的线程(即总共31个将被填充),或者这个2D块线程将被连接,只有最后一个线程将被填充(31*31 = 961; 961%32 = 1) ？
只有一个warp(最后一个)被填充.线程按x,y,z的顺序分组为warp.这样,如果你有一个奇怪的二维数组大小,比如17x17,它连续存储在内存中,你仍然可以从一个17x17线程块中创建32线程warp,这将产生合并的访问.这样,除最后一个之外,所有warp都将生成完全合并的访问.如果在整个过程中使用死线填充单个warp,则在此示例中,在内存访问方面会更浪费.
至少在这个例子中,从机器利用率的角度来看,它的效果更好.
对此的文档支持依赖于理解线程ID和线程索引不相同.
对于给定的主题主题指数由内置变量确定threadIdx.x,threadIdx.y和threadIdx.z.线程ID是唯一的(在线程块内),分配给每个线程的标量数.
线程ID和线程索引之间的关系由此语句给出:
"线程的索引及其线程ID以直接的方式相互关联:对于一维块,它们是相同的;对于二维块大小(Dx,Dy),线程ID是索引(x,y)的线程是(x + y Dx);对于大小为三维的块(Dx,Dy,Dz),索引(x,y,z)的线程的线程ID是(x) + y Dx + z Dx Dy)."
但是线程到warp的分组是由线程ID显式完成的:
"块被分区为warp的方式总是相同的;每个warp包含连续的,增加的线程ID的线程,第一个warp包含线程0."
因此,基于第一个语句,我们看到即使对于像17x17这样的奇数块形状,也没有定义除了在线程块的维度内的线程之外的线程.然后根据第二个语句,通过线程ID连续组合warp会创建warp,所有warp都在其中定义了线程(除了最后一个之外).