是否所有线程在 CUDA 中使用相同数量的寄存器？

Posted 2023-04-15

【中文标题】是否所有线程在 CUDA 中使用相同数量的寄存器？

【英文标题】：Do all threads use the same number of registers in CUDA?

【发布时间】：2021-01-12 20:05:24

【问题描述】：

if (threadIdx.x < 128) 
  float reg[32];
  // do something with reg...
 else 
  return;

假设每个块有 256 个线程，但只有一半的线程在使用寄存器，而另一半在做其他事情（在这种情况下什么都没有）。我的问题是，这个线程块将使用多少个寄存器（仅考虑 reg）？ 32 * 256 还是 32 * 128 ？

【问题讨论】：

除了您的问题（由罗伯特回答）：请注意，数组不会最终成为（线程）本地内存！如果使用非编译时常量索引访问数组，则会发生这种情况

【参考方案1】：

所有线程使用相同数量的寄存器。

这是一个编译时决定，该决定与运行时行为无关，编译器决定网格中所有线程的寄存器使用情况（即内核启动），它不是每个线程决定的。在运行时，必须为每个线程分配必要数量的寄存器，无论它们是否“使用”它们。见here。

您的问题的答案是，无论线程“做什么”，每个块的寄存器数等于每个线程的寄存器数（在编译时确定）乘以每个块的线程数。

所以在您的示例中，它可能是32*256，而不是32*128。

【讨论】：

硬件具有每个 warp 寄存器分配粒度。寄存器分配记录在 CUDA_Occupancy_Calculator.xls GPU 数据表中。对于 CC3.0 - CC8.x，warp 分配粒度为 256 个寄存器/warp。所有线程（无论是活动的还是非活动的）都将分配 8、16、24、32、... 寄存器。