13

在 CUDA 中使用常量的最佳方式是什么？

一种方法是在常量内存中定义常量，例如：

// CUDA global constants
__constant__ int M;

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

另一种方法是使用 C 预处理器：

#define M = ...

我认为用 C 预处理器定义常量要快得多。那么在 CUDA 设备上使用常量内存有哪些好处呢？

4

2 回答 2

19
  1. 在编译时已知的常量应该使用预处理器宏（例如#define）或通过全局/文件范围的 C/C++const变量来定义。
  2. __constant__内存的使用对于使用在内核期间不会改变的某些值并且存在某些访问模式的程序可能是有益的（例如，所有线程同时访问相同的值）。这并不比满足上述第 1 项要求的常数更好或更快。
  3. 如果程序要做出的选择数量相对较少，并且这些选择会影响内核执行，则一种可能的额外编译时优化方法是使用模板化代码/内核
于 2013-11-08T18:41:15.843 回答
8

常规 C/C++ 样式常量：在 CUDA C（本身是 C99 的修改）中，常量是绝对编译时实体。考虑到 GPU 处理的性质，NVCC 中发生的优化量非常重要，这不足为奇。

#define: 宏总是很不优雅，但在紧要关头很有用。

然而，__constant__变量说明符是一种全新的动物，在我看来有点用词不当。我将在下面的空白处记下 Nvidia的内容：

限定词，可选地与__constant__一起使用 __device__，声明一个变量：

  • 驻留在恒定的内存空间中，
  • 具有应用程序的生命周期，
  • 可从网格内的所有线程和主机通过运行时库 (cudaGetSymbolAddress() / cudaGetSymbolSize() / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()) 访问。

Nvidia 的文档指定它__constant__以寄存器级别的速度（接近零延迟）可用，前提是它是一个扭曲的所有线程访问的相同常量。

它们在 CUDA 代码中的全局范围内声明。但是，根据个人（和当前正在进行的）经验，当涉及到单独编译时，您必须小心使用此说明符，例如通过在 C 中放置包装函数将 CUDA 代码（.cu 和 .cuh 文件）与 C/C++ 代码分开样式的标题。

然而，与传统的“常量”指定变量不同，这些变量是在运行时从分配设备内存并最终启动内核的主机代码初始化的。我再说一遍，我目前正在工作的代码演示这些可以在内核执行之前使用 cudaMemcpyToSymbol() 在运行时设置。

考虑到保证访问的 L1 缓存级速度，至少可以说它们非常方便。

于 2014-06-09T03:04:48.073 回答