在CUDA中使用常量

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

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

// CUDA global constants __constant__ int M; int main(void) { ... cudaMemcpyToSymbol("M", &M, sizeof(M)); ... } 

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

 #define M = ... 

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

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

常规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缓存级别速度,至少可以说它们是非常方便的。