为什么指向相同地址的CUDA指针访问比直接访问全局设备内存慢？

性能差异的核心原因是编译器对两种内存访问的优化程度不同，具体可以分为两点： #

• 别名分析限制导致循环不变量无法外提
  你内核的循环逻辑中，ptr[threadIdx.x]是线程维度的固定值，整个循环执行过程中不会发生变化。当直接使用静态__device__变量foo时，编译器明确知道foo是独立于bar的静态全局内存区域，循环内只有对bar的写入操作，没有任何修改foo的逻辑，因此可以直接把foo[threadIdx.x]的读取操作提到循环外部，整个内核执行期间每个线程仅需要做1次全局内存读取，剩下的循环计算全部使用寄存器里的缓存值，执行速度自然极快。
  而ptr是运行时传入的通用指针，编译器没有任何关于它指向的内存区域的先验信息，无法确认ptr和bar是否存在内存重叠（别名问题），为了保证程序执行正确性，编译器不敢把ptr[threadIdx.x]的读取操作外提，只能每次循环都重新从全局内存读取该值。你的循环次数是1<<20即一百万次，相当于每个线程多执行了近一百万次全局内存读取，这是性能出现数倍差异的核心来源。

• 静态符号的额外优化空间
  foo是编译期可知的静态__device__符号，编译器清楚它的地址对齐属性、内存大小，因此可以生成对齐的访存指令，甚至可以利用常量缓存做进一步加速；而通用指针ptr的对齐属性、内存范围都是未知的，编译器只能生成最保守的访存指令，也会带来少量额外开销。

你可以做简单验证：把内核函数的参数修改为const int* __restrict__ ptr，用__restrict__关键字告诉编译器该指针不会和其他内存区域重叠，编译器就可以安全做循环不变量外提优化，此时两种写法的性能会基本一致。

内容的提问来源于stack exchange，提问作者tigertang