我知道对于具有 31 个线程的 1D 线程块，它将被填充到 32 个线程以进行扭曲执行。具有 31*31 线程的 2D 块呢？warp scheduler 是否会为每个维度额外填充 1 个线程（即总共将填充 31 个），或者将这个 2D 块线程连接起来，只填充最后一个线程（31*31=961; 961%32=1） ?
问问题
1375 次
1 回答
9
只有一根经线（最后一根）被填充。线程按 x、y、z 的顺序分组到 warp 中。这样，如果您有一个奇数的 2D 数组大小，例如 17x17，它连续存储在内存中，您仍然可以从 17x17 线程块中创建 32 线程扭曲，从而生成合并访问。通过这种方式，除了最后一个，所有的 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)。”
但是将线程分组为经纱是由线程 ID 明确完成的：
“一个块被划分为warp的方式总是相同的；每个warp包含连续的线程，增加线程ID，第一个warp包含线程0。”
因此，基于第一个陈述，我们看到即使对于像 17x17 这样的奇数块形状，除了线程块维度内的线程之外，没有定义线程。然后基于第二条语句，通过线程 ID 连续组装经纱创建经纱，所有这些经纱都在其中定义了线程（可能最后一个除外。）
于 2013-02-23T19:51:14.387 回答