ein*_*ica -1 c++ cuda constants gpu-warp
nvcc设备代码可以访问内置值,warpSize该值设置为执行内核的设备的warp大小(即在可预见的将来为32).通常你不能把它区分为常数 - 但是如果你试图声明一个长度为warpSize的数组,你就会抱怨它是非常量的...(使用CUDA 7.5)
所以,至少为了这个目的,你有动力去做(编辑):
enum : unsigned int { warp_size = 32 };
在你的标题中的某个地方.但是现在 - 我应该选择哪个,何时？:warpSize,或warp_size？
编辑: warpSize显然是PTX中的编译时常量.问题仍然存在.
tal*_*ies 10
让我们直截了当地说几点.warp大小不是编译时常量,不应该被视为一个.它是一个特定于体系结构的运行时立即数(对于迄今为止所有体系结构,它的值恰好为32).曾几何时,旧的Open64编译器确实向PTX发出了一个常量,但是如果我的记忆没有让我失望,那至少在6年前就会发生变化.
该值可用:
warpSize,where 不是编译时常量(WARP_SZ在这种情况下,编译器会发出PTX 变量).WARP_SZ,它是运行时立即数不要为warp大小声明你自己的常量,这只是要求麻烦.内核数组的正常用例是使用动态分配的共享内存,其大小为warp大小的某个倍数.您可以在运行时从主机API读取warp大小以获取它.如果你有一个静态声明的内核,你需要从warp大小维度,使用模板并在运行时选择正确的实例.后者可能看起来像是不必要的剧院,但对于在实践中几乎从不出现的用例来说,这是正确的做法.这是你的选择.
与塔隆米斯的回答相反，我发现warp_size常数完全可以接受。使用的唯一原因warpSize是使代码与未来可能具有不同大小扭曲的硬件向前兼容。然而，当此类硬件到达时，内核代码很可能还需要进行其他更改才能保持高效。CUDA 不是一种与硬件无关的语言 - 相反，它仍然是一种相当低级的编程语言。生产代码使用随时间变化的各种内部函数（例如__umul24）。
当我们获得不同的扭曲尺寸（例如 64）时，许多事情都会改变：
warpSize必须进行调整int __ballot，虽然int不需要是 32 位，但最常见的是！迭代操作（例如扭曲级别减少）将需要调整其迭代次数。我从未见过有人写：
for (int i = 0; i < log2(warpSize); ++i) ...
对于通常对时间要求严格的代码来说，这会过于复杂。
warpIdx并且laneIdx计算 outthreadIdx需要调整。目前，我看到的最典型的代码是：
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 代码的其余部分可以使用它，而无需任何语法开销。当您决定添加对新架构的支持时，您只需要更改这一段代码即可。