在CUDA中使用常量
哪个是在CUDA中使用常量的最佳方法？
一种方法是在常量内存中定义常量，如：
// CUDA global constants __constant__ int M; int main(void) { ... cudaMemcpyToSymbol("M", &M, sizeof(M)); ... } 一种替代方法是使用C预处理器：
#define M = ...
我认为使用C预处理器定义常量要快得多。 那么在CUDA设备上使用常量内存的好处是什么？
- 编译时已知的常量应使用预处理器宏（例如
#define）或通过全局/文件范围的C / C ++const变量定义。 - 使用
__constant__内存对于使用某些值的程序可能是有益的，这些值在内核的持续时间内不会改变并且存在某些访问模式（例如，所有线程同时访问相同的值）。 这并不比满足上述第1项要求的常数更好或更快。 - 如果程序选择的数量相对较少，并且这些选择会影响内核执行，那么额外编译时优化的一种可能方法是使用模板化代码/内核
常规C / C ++样式常量：在CUDA C（本身是C99的修改）中，常量是绝对编译时实体。 考虑到GPU处理的性质，鉴于NVCC中发生的优化量非常复杂，这并不令人惊讶。
#define ：宏总是非常不优雅，但却很有用。
然而， __constant__变量说明符是一种全新的动物，在我看来是一种误称。 我将在下面的空白处放下Nvidia的内容：
__constant__限定符（可选地与__device__一起使用）声明一个变量：
- 驻留在恒定的存储空间中，
- 有应用程序的生命周期，
- 可以从网格中的所有线程以及主机通过运行时库（cudaGetSymbolAddress（）/ cudaGetSymbolSize（）/ cudaMemcpyToSymbol（）/ cudaMemcpyFromSymbol（））访问。
Nvidia的文档指定__constant__在寄存器级别速度（接近零延迟）时可用，前提是它是由warp的所有线程访问的相同常量。
它们在CUDA代码中声明为全局范围。 然而，基于个人（以及当前正在进行的）经验，在单独编译时必须小心使用此说明符，例如通过在C中放置包装函数将CUDA代码（.cu和.cuh文件）与C / C ++代码分离。风格的标题。
与传统的“常量”指定变量不同，它们在运行时从主机代码初始化，主机代码分配设备内存并最终启动内核。 我再说一遍，我目前正在使用代码来演示这些代码可以在运行时使用cudaMemcpyToSymbol（）在内核执行之前设置 。
考虑到保证访问的L1缓存级别速度，至少可以说它们是非常方便的。