nvcc：结合外部和常量

Posted 2023-02-16

【中文标题】nvcc：结合外部和常量

【英文标题】：nvcc: Combine extern and constant

【发布时间】：2011-10-31 20:11:07

【问题描述】：

我想将我的 CUDA 代码组织成单独的目标文件，以便在编译结束时链接，就像在 C++ 中一样。为此，我希望能够在头文件中声明指向__constant__ 内存的外部指针，并将定义放入其中一个 .cu 文件中，也遵循 C++ 的模式。但似乎当我这样做时，nvcc 忽略了“extern”——它将每个声明作为一个定义。有没有办法解决这个问题？

为了更具体地说明代码和错误，我在头文件中有这个：

extern __device__ void* device_function_table[];

随后在 .cu 文件中：

void* __device__ device_function_table[200];

编译时出现此错误：

(path).cu:40: error: redefinition of ‘void* device_function_table [200]’
(path).hh:29: error: ‘void* device_function_table [200]’ previously declared here

我目前的解决方案是使用 Makefile 魔法将我所有的 .cu 文件组合在一起，并且实际上拥有一个大的翻译单元，但文件组织有些相似。但这已经明显减慢了编译速度，因为对我的任何一个类的更改都意味着重新编译所有类；我预计会增加更多的课程。

编辑：我看到我在文本中输入了__constant__，在示例中输入了__device__；这个问题适用于两者。

【问题讨论】：

CUDA 不支持单独编译，所以我不希望 nvcc 以任何合理的方式处理extern __device__。

【参考方案1】：

长话短说，使用最新的 CUDA 工具包（我在 v8 上）和至少 2.0 的计算能力，在 Visual Studio 中，转到 Project Properties -> CUDA C/C++ -> Common ，找到“Generate Relocatable Device代码”，设置为“是（-rdc=true）”。

对于命令行this page 建议–dc 编译器选项

【参考方案2】：

来自CUDA C Programming Guide 4.0 版，D.2.1.1 节：

不允许使用 __device__ 、 __shared__ 和 __constant__ 限定符 上：

类、结构和联合数据成员， 形式参数， 在主机上执行的函数中的局部变量。

__shared__ 和 __constant__ 变量具有隐含的静态存储。

__device__ 和 __constant__ 变量只能在文件范围内使用。

__device__、__shared__ 和 __constant__ 变量不能使用 extern 关键字定义为外部变量。唯一的例外是动态分配的 __shared__ 变量，如第 B.2.3 节所述。

【讨论】：

谢谢，这回答了我的问题，即使这意味着我不能做我想做的事。

@RolfAndreassen，现在可以了。请参阅下面的答案。

@Tom 也许是时候更新您的答案以包括单独编译了？

【参考方案3】：

从 CUDA 5.0 开始，如果启用了单独编译和链接，现在可以使用 CUDA 获得外部定义的数据。这篇博文解释了它：http://devblogs.nvidia.com/parallelforall/separate-compilation-linking-cuda-device-code/

如果这样做了，人们只需像在原始帖子中一样使用它，它“就可以工作”。