CUDA 共享内存地址空间与全局内存

Question

为了避免我调用的函数非常冗长且不连贯 来自内核的许多设备函数。我分配一个共享 内核调用开始时的缓冲区（每个线程块） 并将指向它的指针传递给所有 device 函数 在内核中执行一些处理步骤。

我想知道以下几点：

如果我在 全局 函数中分配共享内存缓冲区 我通过指针传递的其他 device 函数如何区分 在可能的地址类型（全局设备或共享内存）之间 指针可以引用。

注意用shared修饰符修饰形参是无效的 根据'CUDA编程指南'。 imhoit 的唯一方法是 实现的是

a) 通过在分配的内存上放置标记 b) 通过调用传递不可见的参数。 c) 有一个虚拟的统一地址空间，它有单独的段用于 可以使用全局和共享内存以及对指针的阈值检查吗？

所以我的问题是：我是否需要担心或者应该如何进行 没有将所有函数内联到主内核中？

================================================ =============================================

另一方面，我今天很震惊，带有 CUDA Toolkit 3.0 的 NVCC 不允许所谓的 '来自全局函数的外部调用'，要求它们被内联。这意味着实际上 我必须声明所有 device 函数内联和标题/源的分离 文件损坏。这当然很丑，但是有没有其他的选择？

Answer 1

如果我在全局函数中分配共享内存缓冲区，我传递指针的其他设备函数如何区分指针可能引用的可能地址类型（全局设备或共享内存）。

请注意，在 CUDA 的上下文中，“共享”内存特指块中所有线程之间共享的片上内存。因此，如果您的意思是使用 __shared__ 限定符声明的数组，则使用它在设备函数之间传递信息通常没有意义（因为所有线程都看到相同的内存）。我认为编译器可能将常规数组放入共享内存中？或者它可能在注册文件中。无论如何，它很有可能最终进入全局内存，这将是在设备功能之间传递信息的一种低效方式（尤其是在

另一方面，我今天感到震惊的是，带有 CUDA Toolkit 3.0 的 NVCC 不允许所谓的“来自全局函数的外部调用”，需要内联它们。这意味着实际上我必须内联声明所有设备功能，并且头文件/源文件的分离被破坏。这当然很丑，但是有没有其他的选择？

CUDA 不包含设备代码的链接器，因此您必须将内核和所有相关设备函数保存在同一个 .cu 文件中。

Answer 2

这取决于您的 CUDA 设备的计算能力。对于计算能力= 2.0 的设备，这不是必需的。

默认情况下，内核中的所有函数调用都是内联的，然后编译器在大多数情况下可以使用流分析来查看某些内容是共享的还是全局的。如果您正在为计算能力 warning : Cannot tell what pointer points to, assuming global memory space。这就是编译器无法正确遵循您的指针时得到的结果。