关于 CUDA 中的 const 指针和参数传递答案

【问题标题】：About const pointer and parameter-passing in CUDA关于 CUDA 中的 const 指针和参数传递
【发布时间】：2013-04-15 20:59:49
【问题描述】：

nvcc 如何处理内核中的 const 指针？

根据nvidia的说法，在参数传递过程中为指针添加const和restrict可以使NVCC进行积极的优化，这是严格遵循C/C++的方式吗？

假设 A 是指向数据缓冲区的指针，可能会被其他线程/流频繁更新，但在此测试内核调用期间内容不会被修改：

test<<<blocks, threads>>>(const int *__restrict__ A, int *__restrict__ B);

那么 NVCC 是否可以保持这样的正确性：在每次内核调用时将更新的数据加载到 A 中，而不是加载一些预先缓存的过时数据？

【问题讨论】：

C++ 甚至没有 restrict 关键字，因此它遵循 C99 而不是 C++。但一般来说，潜在的混叠会扼杀优化。
另外，如果数据是从其他线程更新的，您不应该将其设为volatile 或以其他方式包含内存屏障吗？这在 CPU 端 C 和 C++ 中肯定是必需的。
@BenVoigt：数据由其他线程修改，但不会在此内核调用期间和期间修改，我认为我在这里不需要 volatile（在我的情况下，有一个全局标志所有线程来控制这个，但那是 CPU 方面的东西），至于限制，它是一个 NVCC 认可的关键字。
const 表示“在此范围内将此对象视为只读”。编译器通常可以检测与此修饰符冲突的用法。限制意味着“此指针是在此范围内访问指向对象的唯一路径”。这是程序员做出的断言，编译器通常无法验证。如果程序员违背了他们的承诺（即实际上存在别名），则代码必然不会按照程序员的意愿工作。由于 C++（即使是 2011 年的版本）不支持受限指针，CUDA 遵循 C99，但使用关键字 restrict。

标签： c++ c cuda nvidia

【解决方案1】：

const 像 C++ 一样工作。 const 变量无法更改，编译器会在编译时对其进行检查。编译器仅检查给定范围的 const 正确性，因为可以使用 C 样式转换来更改 constness。

restrict 以 C 方式工作。当您将指针标记为限制时，编译器假定这些指针没有别名。这是你给定的事实，编译器不会检查这个事实是否正确。

关于您的问题，NVCC 不会确保内核启动之间的全局内存写入和读取的正确性。由于内核启动在 CUDA 中是异步的，因此您必须确保修改这些内存空间的内核不会同时执行。您可以通过同步内存副本和/或cudaDeviceSynchronize() 来实现这一点。如果同时启动这些内核，则无法确保在来自其他内核的访问之前，不同内核的所有更改都已提交到全局内存。

【讨论】：