在 PTX 中使用共享内存时出现“意外的地址空间”编译错误答案

【问题标题】："Unexpected address space" compilation error while using shared memory in PTX在 PTX 中使用共享内存时出现“意外的地址空间”编译错误
【发布时间】：2013-03-07 06:55:00
【问题描述】：

我编写了一个简单的内核，在其中我将共享内存数组声明为

extern __shared__ float As[100];

在我的内核启动中，我指定了 number_of_bytes 的共享内存。编译内核（到 PTX）时出现错误“意外的地址空间”。我正在使用来自 svn 的相当新版本的 LLVM（3.3 正在进行中）。有什么想法我在这里做错了吗？问题似乎与 extern 关键字有关，但是我还要如何指定它呢？（共享内存）。我应该使用不同的 LLVM 构建吗？

配置 CUDA 5.0 , Nvidia Tesla C1060

【问题讨论】：

试试这个：extern __shared__ float As[];
所以您使用的是自己的 LLVM 构建，而不是 CUDA 工具包中的任何东西？
也许你可以提供更完整的源代码。
您的代码是否使用默认的 CUDA 5 工具包正确构建？
@talonmies 不是我自己的，而是来自大约一个月前拍摄的 LLVM svn。 Robert Crovella：好吧，nvcc 会抛出“错误：local 和 shared 变量不能有外部链接”

标签： cuda llvm ptx

【解决方案1】：

好吧，根据 Nvidia 论坛的 Gert-Jan，在这种情况下实际上并不需要 extern 关键字。我不确定他在 SO 上的 id 是什么。

他的回复——

"如果你知道你的共享内存数组有多少个元素（例如 100 个元素），你不应该使用 extern 关键字，并且你不必在内核启动时指定共享内存的字节数（编译器可以自己算出来）。只有当你不知道你需要多少元素时，你必须在内核启动时指定这个，并且在你的内核中你必须写“extern shared 浮动 *As"。"

希望这对其他用户有所帮助。

【讨论】：

如果您事先知道需要多少空间，则确实不需要extern。但它也应该是有效的语法，说extern __shared__ float As[] 并在内核调用时指定大小。所以对我来说，这个答案什么也没解释。 sgar91 建议你修复你的 extern 语法，你说你试过了，得到了同样的错误。据了解，指定共享内存的有效语法有两种。两者都应该工作。您尝试了一个，它不起作用，然后您切换到另一个，它确实起作用了。这并不能解释为什么第一种语法不起作用。

【解决方案2】：

我不确定 CUDA-C/C++ 是否支持此功能，但可能尝试将地址空间属性设置为解决方法：

__attribute__((address_space(3)))
extern __shared__ float As[100];

那应该强制 llvm 把它放在共享地址空间.... 祝你好运！

【讨论】：