如何以程序员愉快的方式使用 CUDA 常量内存？答案

【问题标题】：How to use CUDA constant memory in a programmer pleasant way?如何以程序员愉快的方式使用 CUDA 常量内存？
【发布时间】：2011-04-29 19:29:26
【问题描述】：

我正在使用 CUDA 框架开发一个数字运算应用程序。我有一些所有线程都应该可以访问的静态数据，所以我把它放在常量内存中，如下所示：

__device__ __constant__ CaseParams deviceCaseParams;

我使用调用 cudaMemcpyToSymbol 将这些参数从主机传输到设备：

void copyMetaData(CaseParams* caseParams)
{
    cudaMemcpyToSymbol("deviceCaseParams", caseParams, sizeof(CaseParams));
}

哪个有效。

无论如何，似乎（通过反复试验，以及阅读网上的帖子）由于某种病态的原因，deviceCaseParams 的声明和它的复制操作（对 cudaMemcpyToSymbol 的调用）必须在 同一个文件。目前我在 .cu 文件中有这两个，但我真的希望在 .cuh 文件中有参数结构，以便任何实现都可以看到它。这意味着我还必须在 a 头文件中有 copyMetaData 函数，但这会破坏链接（符号已定义），因为 .cpp 和 .cu 文件都包含此头文件（因此 MS C++ 编译器和 nvcc 都会编译它)。

这里有人对设计有什么建议吗？

更新：查看 cmets

【问题讨论】：

您确定它们必须在同一个文件中，而不仅仅是在同一个翻译单元中吗？（即声明可以在头文件中，然后将其#included 到源文件中）。
几分钟前我试过了，看来你是对的。不过，我不明白上次尝试时出了什么问题。它现在肯定有效。谢谢。

标签： c++ visual-studio header linker cuda

【解决方案1】：

如果您在运行时查找符号（即通过将字符串作为第一个参数传递给cudaMemcpyToSymbol 就像你在你的例子中一样）。

此外，对于 Fermi 类设备，您只需 malloc 内存 (cudaMalloc)，复制到设备内存，然后将参数作为 const 指针传递。编译器将识别您是否正在跨经线统一访问数据，如果是，将使用常量缓存。有关详细信息，请参阅 CUDA 编程指南。注意：您需要使用-arch=sm_20 进行编译。

【讨论】：

【解决方案2】：

如果您使用的是 pre-Fermi CUDA，您现在会发现这个问题不仅适用于常量内存，它适用于您想要的 CUDA 方面的任何事情。我发现的唯一两种方法是：

将所有 CUDA 写入单个文件 (.cu)，或
如果您需要将代码拆分为单独的文件，请将自己限制为您的单个 .cu 文件随后包含的标头。

如果您需要在 CUDA 和 C/C++ 之间共享代码，或者在项目之间共享一些公共代码，选项 2 是唯一的选择。开始看起来很不自然，但它解决了问题。您仍然可以构建您的代码，只是不是以典型的 C 类方式。主要开销是每次构建时都会编译所有内容。这样做的好处（我认为这可能是它以这种方式工作的原因）是 CUDA 编译器可以一次访问所有源代码，这有利于优化。

【讨论】：