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

【问题标题】：How to define a CUDA shared memory with a size known at run time?如何定义运行时已知大小的 CUDA 共享内存？
【发布时间】：2012-04-13 17:58:45
【问题描述】：

CUDA 中的__shared__ 内存似乎在编译时需要一个已知的大小。但是，在我的问题中，__shared__ 内存大小仅在运行时知道，即

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

这会以“错误：常量值未知”结束，我不知道如何解决这个问题。

【问题讨论】：

【解决方案1】：

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

但是，支持动态指定单个 __shared__ 数组大小的特殊情况，该数组对于所有线程的大小都相同。见allocating shared memory。

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

【讨论】：

实际上我试图将 blockDim.x 设置为 shared 内存大小，这对于同一块中的所有线程都是相同的，但它仍然失败（与一个不同的错误）
@HailiangZhang：您可能不想计划通过动态改变块尺寸来解决您的问题。通常，您可以通过仔细考虑内核的资源使用情况并在此基础上设置固定的最佳块尺寸来获得最佳性能。例如，使用 CUDA 占用计算器电子表格。此外，为了获得最佳性能，尺寸应乘以扭曲尺寸的倍数。您只能动态调整网格尺寸，以适应您的数据。
@RogerDahl：这个答案不正确。您可以在运行时动态确定内核共享内存 - 这是自 1.0 以来 CUDA 的一项功能。具体方法请参见this answer。
@talonmies：啊，所以这就是那个神秘的第三个参数是为了什么！生活和学习...谢谢。我已经确定了答案。

【解决方案2】：

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

【讨论】：