在CUDA中使用常量

Question

哪个是在CUDA中使用常量的最佳方法？

一种方法是在常量内存中定义常量,例如:

// CUDA global constants
__constant__ int M;

int main(void)
{
    ...
    cudaMemcpyToSymbol("M", &M, sizeof(M));
    ...
}

一种替代方法是使用C预处理器:

#define M = ... 

我认为使用C预处理器定义常量要快得多.那么在CUDA设备上使用常量内存的好处是什么？

Answer 1

1. 编译时已知的常量应使用预处理器宏(例如#define)或通过const全局/文件范围的C/C++ 变量定义.
2. 的使用__constant__内存可能是谁使用不为内核的时间发生变化,其某些访问模式存在(例如,所有线程访问在同一时间相同的值)的特定值的程序是有益的.这并不比满足上述第1项要求的常数更好或更快.
3. 如果程序选择的数量相对较少,并且这些选择会影响内核执行,那么额外编译时优化的一种可能方法是使用模板化代码/内核

Answer 2

常规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高速缓存级别的速度，它们非常方便。