如何在运行时定义大小已知的CUDA共享内存？

Question

在CUDA中，__shared__内存似乎需要在编译时知道其大小。然而，在我的问题中，__shared__内存的大小只有在运行时才能确定，即：

int size=get_size();
__shared__ mem[size];

这将导致 "错误：常量值未知"，我不确定如何解决这个问题。

- Hailiang Zhang

---

可能是分配共享内存的重复问题。 - talonmies

Answer 1

共享内存的目的是允许块中的线程协作。当您将一个数组声明为__shared__时，块中的每个线程看到的是相同的内存，因此对于给定的线程来说，为共享内存中的数组设置自己的大小是没有意义的。

然而，支持动态指定单个__shared__数组的大小，该大小对于所有线程都相同。请参见分配共享内存。

如果确实需要为每个线程动态分配内存，则可以在内核中使用new或malloc（在Fermi上），但它们分配的是全局内存，可能会很慢。

- Roger Dahl

---

实际上，我试图将blockDim.x作为__shared__内存大小，这对于同一块中的所有线程都是相同的，但它仍然失败了（尽管出现了不同的错误）。 - Hailiang Zhang

@HailiangZhang：你可能不想通过动态变化块维度来解决问题。通常情况下，你应该仔细考虑内核的资源使用情况，并基于此设置一个固定的、最优的块维度，以获得最佳性能。例如，可以使用CUDA占用率计算器电子表格。此外，维度应该乘以warp大小的倍数以获得最佳性能。你只需要动态调整网格维度以适应你的数据。 - Roger Dahl

@RogerDahl：这个答案是错误的。你__可以__在运行时动态确定内核共享内存 - 这是CUDA自1.0以来的一个特性。请参见此答案了解详情。 - talonmies

@talonmies：啊，原来那个神秘的第三个参数是用来干嘛的！学到了...感谢您。我已经修正了答案。 - Roger Dahl

Answer 2

你应该使用extern__shared__ mem[];（动态共享内存）而不是__shared__ mem[size];（静态共享内存）。参见[https://devblogs.nvidia.com/parallelforall/using-shared-memory-cuda-cc/][1]]。

- Hosein Hasani