在 Cuda 内核的嵌套数组中混合使用静态和动态共享内存

Question

我有时会在 CUDA 内核中看到以下共享内存声明，但我不确定它的含义：

extern __shared__ T shmem[][SZ]

作为SZ编译时常量。内核启动如下：

kernel<<<grid, block, shared_memory_size>>>()

我的问题是：

• 为什么动态和静态共享内存的混合会很有用（除了简化程序员的地址计算之外）？
• 共享内存缓冲区 shmem 的总大小是多少？是吗shared_memory_size * SZ？
• 与问题2相关：假设我可以在编译时计算shared_memory_size，我将如何重写共享内存声明以使其成为静态？shared T shmem[SZ*shared_memory_size]？

Answer 1

__extern__ shared T shmem[][SZ]

正确的语法是

extern __shared__ T shmem[][SZ];

为什么动态和静态共享内存的混合会很有用（除了简化程序员的地址计算之外）？

我相信地址计算是唯一的原因。当然，这也影响静态分析等，但仅此而已；让您的工作更轻松。

当然，如果您创建SZ一个运行时变量，即使它可能是编译时的，那么计算成本也会上升，因为 GPU 必须为其地址生成执行适当的整数乘法，而不是将其优化为更便宜的操作（例如位移位）。

共享内存缓冲区 shmem 的总大小是多少？是吗shared_memory_size * SZ？

不，内核启动参数以字节为单位。所以沿外部尺寸的尺寸将为shared_memory_size / (SZ * sizeof(T))。换句话说，如果你想shmem[N][SZ]在运行时，那么shared_memory_size = N * SZ * sizeof(T)

假设我可以在编译时计算shared_memory_size如何重写共享内存声明才能使其成为静态？shared T shmem[SZ*shared_memory_size]？

你的意思是这样吗？

__shared__ T shmem[N][SZ];

请注意它是如何不再声明的extern。

脚注：就命名法而言，我不会将其称为静态和动态共享内存的混合。我相信大多数人都会将混合解释为在同一个内核中使用__shared__ T和。extern __shared__ T