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

Question

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

__extern__ shared T shmem[][SZ]

其中

SZ

是编译时常量。内核启动为：

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

我的问题是：

• 为什么动态和静态共享内存的混合会很有用（除了简化程序员的地址计算之外）？
• 共享内存缓冲区shmem的总大小是多少？是

  shared_memory_size * SZ

  吗？
• 与问题 2 相关：假设我可以在编译时计算共享内存大小，我将如何重写共享内存声明以使其成为静态？

  shared T shmem[SZ*shared_memory_size]

  ？

Answer 1

__extern__ shared T shmem[][SZ]

正确的语法是

extern __shared__ T shmem[][SZ];

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

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

共享内存缓冲区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];