我刚刚问过自己，使用例如threadIdx.x进行复杂的索引计算是否会对性能产生影响。内核一上传到设备，这些变量就变为常量了吗？
我想导航到一个巨大的数组，其中索引取决于threadIdx.x、threadIdx.y和threadId*.z。我需要例如之类的模运算
array[threadIdx.y % 2 + ...]
在索引计算中有一个加法和一个模数。
来自CUDA编程指南：operator+的吞吐量非常高（3.5计算能力的GPU为160）。
CCD_ 2需要数十个具有类似于CCD_ 3的吞吐量的操作。
在您的情况下，您使用的是带有文字常量的operator%，编译器很可能会对其进行优化。此外，您的常数是2的幂，因此编译器将用按位的operator&（与operator+的吞吐量相同）替换它。
对应用程序进行概要分析是很重要的，这样可以避免在优化算术运算时浪费时间，而不会获得任何性能。算术运算通常被内存加载和存储操作完全隐藏，在这种情况下，您需要专注于优化内存吞吐量。
我假设
array[threadIdx.y % 2 + ...]
只是一个例子。
一般来说，%操作可能较慢。加快指数计算的一个有用技巧是注意
foo%n==foo&(n-1) if n is a power of 2
因此，对于上面的例子，编译器可能会为您进行优化，但如果您有foo%n，那么上面的技巧值得使用。
如果有人感兴趣，我已经评估了相应的PTX代码。
（1） 复杂的线程ID计算会对性能产生影响。"threadIdx.x"等不是常数。
（2） "threadIdx.y%2"是有效实现的，对应于"threadIdx.y&0x00000001"（Cuda Toolkit 5.5）。