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

Question

给出了一个如何在 CUDA 中使用动态分配和外部 shared 内存的示例：Use dynamic shared memory allocation for two different vectors

    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

使用动态分配的共享内存（而不是静态分配）的原因类似于您可能希望动态分配 anything 而不是静态分配的原因之一：在编译时，您不需要'不知道你想要分配的大小。

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

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

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

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

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

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

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

Answer 2

使用共享内存的另一个正当理由是内存合并，否则可能是串行操作。参考http://devblogs.nvidia.com/parallelforall/using-shared-memory-cuda-cc/