我什么时候应该使用 CUDA 的 built-in warpSize，而不是我自己的常量？

Question

nvcc 设备代码可以访问 built-in 值，warpSize，该值设置为执行内核的设备的 warp 大小（即在可预见的未来为 32）。通常你不能将它与常量区分开来——但是如果你试图声明一个长度为 warpSize 的数组，你会收到一个关于它 non-const...（使用 CUDA 7.5）

的投诉

因此，至少出于这个目的，您有动力拥有类似 (edit):

enum : unsigned int { warp_size  = 32 };

在您的 headers 中的某处。但是现在 - 我应该更喜欢哪个，什么时候？：warpSize，还是warp_size？

编辑： warpSize 显然是 PTX 中的一个 compile-time 常量。尽管如此，问题依然存在。

Answer 1

让我们直截了当地说两点。扭曲大小不是编译时间常数，不应将其视为一个。它是特定于体系结构的运行时即时常量（对于迄今为止的所有体系结构，它的值恰好是 32）。曾几何时，旧的 Open64 编译器确实向 PTX 发出了一个常量，但是如果我的记忆没有让我失望的话，那至少在 6 年前就发生了变化。

可用值：

在 CUDA C 中通过 warpSize，其中 不是编译时间常量 （在这种情况下，编译器会发出 PTX WARP_SZ 变量).
在 PTX 汇编器中通过 WARP_SZ，它是一个运行时立即常量
从运行时API作为设备property

不要声明你自己的 warp 大小常量，那只会自找麻烦。尺寸为 warp 大小的某个倍数的内核数组的正常用例是使用动态分配的共享内存。您可以在运行时从主机 API 读取 warp 大小来获取它。如果你有一个静态声明的内核，你需要从 warp 大小确定尺寸，在运行时使用模板和 select 正确的实例。后者可能看起来像是不必要的戏剧，但对于在实践中几乎从未出现过的用例来说，这是正确的做法。选择权在你。

Answer 2

与 talonmies 的回答相反，我发现 warp_size 常量完全可以接受。使用 warpSize 的唯一原因是使代码 forward-compatibly 与未来可能具有不同大小扭曲的硬件一起使用。然而，当这样的硬件到来时，内核代码很可能还需要进行其他更改以保持高效。 CUDA 不是 hardware-agnostic 语言 - 相反，它仍然是一种 low-level 编程语言。生产代码使用各种随时间变化的内在函数（例如 __umul24）。

当我们获得不同的 warp 大小（例如 64）的那一天，很多事情都会改变：

warpSize显然要调整
许多 warp-level 内在函数需要调整其签名，或生成新版本，例如int __ballot，虽然 int 不需要是 32 位的，但是最常见的！
迭代操作，例如 warp-level 缩减，需要调整迭代次数。没见过有人写：
```
for (int i = 0; i < log2(warpSize); ++i) ...
```
在通常是 time-critical 一段代码的情况下，这会过于复杂。
warpIdx 和 laneIdx 计算 threadIdx 需要调整。目前，我看到的最典型的代码是：
```
warpIdx = threadIdx.x/32;
laneIdx = threadIdx.x%32;
```
简化为简单的 right-shift 和掩码操作。但是，如果将 32 替换为 warpSize，这会突然变成一个非常昂贵的操作！

同时，在代码中使用 warpSize 会阻止优化，因为正式来说它不是 compile-time 已知常量。此外，如果共享内存的数量取决于 warpSize，这会强制您使用动态分配的 shmem（根据 talonmies 的回答）。然而，它的语法使用起来很不方便，尤其是当你有多个数组时——这迫使你自己进行指针运算并手动计算所有内存使用量的总和。

为此使用模板 warp_size 是一个部分解决方案，但在每次函数调用时增加了一层语法复杂性：

deviceFunction<warp_size>(params)

这会混淆代码。样板文件越多，代码越难阅读和维护。

我的建议是使用一个 header 来控制所有 model-specific 常量，例如

#if __CUDA_ARCH__ <= 600
//all devices of compute capability <= 6.0
static const int warp_size = 32; 
#endif

现在您的其余 CUDA 代码可以使用它而无需任何语法开销。当您决定添加对更新架构的支持时，您只需要更改这一段代码。

我什么时候应该使用 CUDA 的 built-in warpSize，而不是我自己的常量？

When should I use CUDA's built-in warpSize, as opposed to my own proper constant?

c++

cuda

constants

gpu-warp