我什么时候应该使用 CUDA 的内置 warpSize，而不是我自己的适当常量？答案

【问题标题】：When should I use CUDA's built-in warpSize, as opposed to my own proper constant?我什么时候应该使用 CUDA 的内置 warpSize，而不是我自己的适当常量？
【发布时间】：2016-07-03 00:04:46
【问题描述】：

nvcc 设备代码可以访问内置值warpSize，该值设置为执行内核的设备的扭曲大小（即在可预见的将来为 32）。通常你不能把它和一个常量区分开来——但是如果你尝试声明一个长度为 warpSize 的数组，你会收到一个关于它是非常量的抱怨......（使用 CUDA 7.5）

所以，至少出于这个目的，你有动力拥有类似 (edit) 的东西：

enum : unsigned int { warp_size  = 32 };

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

编辑： warpSize 显然是 PTX 中的编译时常量。不过，问题仍然存在。

【问题讨论】：

您所说的“来自...来自指令本身”是什么意思？
您的很多问题都可以通过阅读一点文档来回答 - docs.nvidia.com/cuda/parallel-thread-execution/…
@talonmies：有时你不知道要阅读哪个文件以及在哪里阅读。
@njuffa：我的意思是来自机器指令（SASS 指令）的操作数。

标签： c++ cuda constants gpu-warp

【解决方案1】：

让我们直说两点。经线大小不是编译时间常数，不应被视为一个。它是一个特定于架构的运行时立即常量（对于迄今为止的所有架构，它的值恰好是 32）。曾几何时，旧的 Open64 编译器确实向 PTX 发出了一个常量，但是如果我没有记错的话，至少在 6 年前就发生了变化。

值可用：

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

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

【讨论】：

这个答案是正确的/最好的/优雅的。然而，有很多代码将扭曲大小视为一个通用常数（有些甚至起源于 NVIDIA，因此遵循关于住在玻璃房子里的人的规则......）。因此，一个可能的折衷方案是将 warpsize 视为一个通用常量，但在代码开头使用断言或类似方法测试硬件返回的实际运行时立即值是否等于假定的通用常量.这应该是安全且足够的，但需要注意的是，如果它发生变化，您的代码将会中断。
@RobertCrovella：“代码的开头”

【解决方案2】：

与 talonmies 的回答相反，我发现 warp_size 常量完全可以接受。使用warpSize 的唯一原因是使代码与未来可能具有不同大小扭曲的硬件向前兼容。然而，当这样的硬件到来时，内核代码很可能还需要其他更改才能保持效率。 CUDA 不是一种与硬件无关的语言——相反，它仍然是一种相当低级的编程语言。生产代码使用随时间变化的各种内在函数（例如__umul24）。

当我们得到不同的经纱尺寸（例如 64）时，很多事情都会改变：

warpSize 显然需要调整
许多经线级别的内在函数需要调整其签名，或生成新版本，例如int __ballot，虽然 int 不需要是 32 位的，但它最常见！
迭代操作，例如扭曲级别的缩减，将需要调整其迭代次数。我从没见过有人写：
```
for (int i = 0; i < log2(warpSize); ++i) ...
```
这对于通常是时间关键的代码段来说过于复杂。
warpIdx 和 laneIdx 的 threadIdx 计算需要调整。目前，我看到的最典型的代码是：
```
warpIdx = threadIdx.x/32;
laneIdx = threadIdx.x%32;
```
这简化为简单的右移和掩码操作。但是，如果您将32 替换为warpSize，这会突然变得非常昂贵！

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

为 warp_size 使用模板是部分解决方案，但会增加每次函数调用所需的语法复杂性：

deviceFunction<warp_size>(params)

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

我的建议是使用一个标头来控制所有特定于模型的常量，例如

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

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

【讨论】：