0

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

所以，至少为了这个目的，你有动力去拥有类似（编辑）的东西：

enum : unsigned int { warp_size  = 32 };

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

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

4

2 回答 2

10

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

该值可用：

  1. 在 CUDA C viawarpSize中，其中 is不是编译时间常数WARP_SZ（在这种情况下，编译器会发出PTX变量）。
  2. 在 PTX 汇编器 viaWARP_SZ中，它是运行时立即常数
  3. 从运行时 API 作为设备属性

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

于 2016-03-17T07:11:13.643 回答
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 代码的其余部分可以使用它而无需任何语法开销。当您决定添加对更新架构的支持时，您只需要更改这段代码。

于 2017-02-21T16:57:10.960 回答