如何以程序员友好的方式使用CUDA常量内存？

Question

我正在使用CUDA框架开发一个进行数字计算的应用程序。我有一些静态数据应该对所有线程都可访问，所以我把它放在常量内存里，像这样：

__device__ __constant__ CaseParams deviceCaseParams;

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

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

这似乎是一个奇怪的问题，通过试错和阅读网络帖子得知，deviceCaseParams声明和它的复制操作(call to cudaMemcpyToSymbol)必须在同一个文件中。目前我把这两个操作都放在了一个.cu文件中，但我真的希望将参数结构体放在一个.cuh文件中，以便任何实现都可以看到它。这意味着我也必须在头文件中包含copyMetaData函数，但这会导致链接出错(符号已定义)，因为.cpp和.cu文件都包含了这个头文件(因此MS C++编译器和nvcc都编译它)。

有人对这里的设计有任何建议吗？

更新:请参见评论。

- Yngve Sneen Lindal

---

你确定它们必须在同一个文件中，而不仅仅是在同一个翻译单元中吗？（即声明可以在头文件中，然后被 #include 到源文件中）。 - Oliver Charlesworth

我几分钟前试了一下，看来你是对的。不过我不明白上次我尝试时到底出了什么问题。现在它绝对可以用了。谢谢。 - Yngve Sneen Lindal

Answer 1

如果您的CUDA版本是最新的（例如3.2），则可以在不同的翻译单元中执行memcpy操作，只需在运行时查找符号（即通过将字符串作为第一个参数传递给cudaMemcpyToSymbol，就像您在示例中所做的那样）。

此外，对于Fermi级设备，您可以直接使用malloc函数分配内存（cudaMalloc），将数据复制到设备内存中，然后将其作为常量指针传递。编译器会识别是否在warp中统一访问数据，如果是，则会使用常量缓存。有关更多信息，请参见CUDA编程指南。注意：您需要使用“-arch = sm_20”编译。

- Tom

Answer 2

如果您使用的是Fermi之前的CUDA版本，您肯定已经发现这个问题不仅适用于常量内存，而且也适用于CUDA的任何东西。我发现解决此问题的唯二方法是要么：

1. 在一个单独的文件中（.cu）编写所有CUDA代码；或者
2. 如果您需要将代码分成多个文件，请仅使用头文件，然后在单个.cu文件中包含这些头文件。

如果您需要在CUDA和C/C++之间共享代码或者在项目之间共享一些公共代码，则选项2是唯一的选择。起初似乎非常不自然，但它能够解决问题。您仍然可以构造您的代码，只是不能像典型的C语言那样。主要的开销是每次构建时都需要编译所有代码。优点是（我认为可能是为什么它以这种方式工作的原因）CUDA编译器可以一次性访问所有源代码，这对于优化是有利的。

- Raffles