当我以 32 宽的 warp CUDA 架构为目标时，我应该使用 warpSize 吗？

Posted 2023-02-22

【中文标题】当我以 32 宽的 warp CUDA 架构为目标时，我应该使用 warpSize 吗？

【英文标题】：When I target 32-wide warp CUDA architectures, should I use warpSize?

【发布时间】：2017-02-21 15:23:25

【问题描述】：

这是this one 的后续问题。

假设我有一个 CUDA 内核

template<unsigned ThreadsPerWarp>
___global__ foo(bar_t* a, const baz_t* b);

我正在针对 ThreadsPerWarp 为 32 岁的情况对其进行专门化（这绕过了对 Talonmies 对我之前问题的回答的有效批评。）

在这个函数（或从它调用的其他__device__ 函数）的主体中——我应该更喜欢使用ThreadsPerWarp 的常量值吗？还是使用warpSize 更好？或者 - 就它生成的 PTX 而言，编译器是否都一样？

【问题讨论】：

从优化的角度来看，使用ThreadsPerWarp 应该更好。

@RobertCrovella：总是无一例外？也就是说，nvcc 没有比32 更清楚地“注意到”warpSize 的优化器？毕竟，talonmies 在链接到的问题中说，生成的 PTX 不假定扭曲大小为 32。

在来自@talonmies 的the answer you linked 中，他说“同时，在代码中使用warpSize 会阻止优化，因为从形式上讲，它不是编译时已知的常量。”我几乎只是重申这一点（我同意，显然）。我猜你现在要求证明某些东西不存在 - 更难做到。

@RobertCrovella：我实际上在想，也许确实存在类似的东西，否则 talonmies 就不会那么强烈地支持warpSize。但你基本上已经用你的 cmets 回答了我的问题。

【参考方案1】：

不，不要使用warpSize。

似乎除了潜在的未来证明（这在实践中是有问题的）之外，使用它没有任何优势。相反，你可以很好地使用类似的东西：

enum : unsigned  warp_size = 32 ;