为什么我应该使用 CUDA __shared__ 内存作为“extern”

Question

给出了如何在 CUDA 中使用动态分配并因此使用外部共享内存的示例：Use dynamic shared memory allocation for two different vector

    extern __shared__ float array[];
    __device__ void func()      // __device__ or __global__ function
    {
        short* array0 = (short*)array; 
        float* array1 = (float*)&array0[128];
        int*   array2 =   (int*)&array1[64];
    }

但是，当我必须手动将其分配给变量时，为什么还要使用 extern 动态分配的共享内存呢？

我认为以下解决方案没有任何缺点：

    __device__ void func()      // __device__ or __global__ function
    {
        __shared__ float array[MAXIMALLY_NEEDED_SIZE];
        short* array0 = (short*)array;
        float* array1 = (float*)&array0[128];
        int*   array2 =   (int*)&array1[64];
    }

显然，使用第一个解决方案，我可以节省一些共享内存。但这对我有什么帮助呢？

（我问动态分配内存有一个很好的理由，但我没有看到，所以我可能缺乏理解。这就是我问的原因。）

Answer 1

使用动态分配的共享内存（而不是静态分配）的原因类似于您可能想要动态而不是静态分配任何东西的原因之一：在编译时，您不知道您想要的分配大小.

你给出的例子并没有很好地说明这一点。该示例的最初目的是说明如何在动态分配的情况下处理驻留在共享内存中的多个独立对象，而不是强调动态与静态共享内存的使用。

显然，使用第一个解决方案，我可以节省一些共享内存。但这对我有什么帮助呢？

保存共享内存可能很有价值的一个可能原因是，它会影响占用率，从而影响性能。

假设我有一个并行归约代码，并假设它使用共享内存作为主要归约媒介。通常，我需要的共享内存量与我在线程块中使用的线程数有关。现在让我们还假设根据我遇到的确切问题，我可能想在运行时调整每个线程块的线程数。

如果我启动一个包含 256 个线程的线程块，并且我正在对 64 位整数进行并行缩减，则每个线程块可能需要 256*8 字节 (2KB) 的共享内存。如果我启动一个包含 1024 个线程的线程块，我将需要每个线程块 8KB 的共享内存（这是可以想象的最大值）。

如果我只是硬编码这个值，以便它可以在编译时作为静态分配的一部分使用，我将需要使用 8KB 值。这将限制我在大多数 GPU 上最多占用 6 个线程块（6*8KB = 48KB 最大共享内存），即使我启动的线程块只有 256 个线程。（如果我出于任何其他目的需要任何共享内存，那么我的最大占用将少于 6 个线程块。）

通过动态分配，1024 个线程的线程块仍然具有与上述相同的限制，但是使用 256 个线程启动的线程块将能够实现理论上更高的占用率（至少基于共享内存限制），这可以转化为更高的性能。