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

Question

一个示例展示了如何在CUDA中使用动态分配的、因此是extern的__shared__内存：使用动态共享内存分配两个不同向量。

    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];
    }

显然，使用第一种解决方案，我可以节省一些共享内存。但这又如何有助于我呢？

（我猜动态分配内存有一个很好的原因，但我没看出来，所以我可能在理解上有欠缺。这就是为什么我会问的原因。）

- Michael

---

在第一个代码示例中，“extern”关键字意味着“array”是指向其他地方声明的共享内存。在第二个代码示例中，缺少“extern”关键字意味着所有引用“array”的共享内存都在那一点被声明。我相信你必须使用两个声明：在一个地方你不使用“extern”，在所有其他使用该共享内存的地方你都需要使用。 - Evil Dog Pie

此外，有关为什么动态内存始终声明为 extern 的信息，请参见此链接。简而言之，这是因为在启动内核时将内存分配给线程块。也就是说，它对于您的所有代码都是 extern，因为它位于内核分配的线程块共享内存中。 - Evil Dog Pie

Answer 1

使用动态分配的共享内存（而不是静态分配）的原因类似于您可能希望动态分配任何东西而不是静态分配的原因：在编译时，您不知道要分配的大小。您提供的示例并没有很好地说明这一点。该示例最初的目的是说明如何处理驻留在动态分配情况下的共享内存中的多个独立对象，而不是突出使用动态与静态共享内存的用途。

引用：“显然，通过第一个解决方案，我可以节省一些共享内存。但这怎么会对我有所帮助呢？”

节省共享内存可能有价值的一个可能原因是它可能会影响占用率，从而影响性能。假设我有一个并行约简代码，并且假设它使用共享内存作为主要约简介质。通常，我需要的共享内存量将与我在线程块中使用的线程数相关。现在让我们还假设，根据我实际遇到的问题，我可能希望在运行时调整每个线程块中的线程数。

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

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

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

- Robert Crovella

---

"我可能想在运行时调整每个线程块的线程数。" <- 好的，这是一个完美的一句话回答问题，谢谢！ :-) - Michael

Answer 2

另一个使用案例是代码分析/基准测试，因为网格和动态分配的内存可以根据运行时的用户提供的参数决定，而无需昂贵的重新编译。

考虑一个动态分配共享内存的内核。

__global__ void kernel(...) {
    extern __shared__ int sdata[];
    ...

由...推出

kernel<<<gridSize, blockSize, sharedBytes>>>(...)

其中gridSize, blockSize, sharedBytes是参数，可以作为输入参数提供。

- Maciej Skorski

Answer 3

为了避免像 float* array1 = (float*)&array0[128] 这样混乱的声明，您可以将共享内容打包在一个结构体中：

struct Shared
{
    short array0[128];
    float array1[64];
    int array2[32];
};

__device__ void func()
{
    extern __shared__ Shared shr[];
    /* use example:
       shr->array0[some_index]
       shr->array1[some_index]
       shr->array2[some_index]
    */
}

- user2052436