memory - CUDA 共享内存和全局内存有什么区别？

Question

我对如何在 CUDA 中使用共享内存和全局内存感到困惑，尤其是在以下方面:

• 当我们使用 cudaMalloc() 时，我们得到一个指向 shared 还是 global 的指针内存力？
• 全局内存是驻留在主机还是设备上？
• 有没有大小限制为任一？
• 哪个访问速度更快？

• 正在存储一个共享内存中的变量与通过核心？ IE。而不是拥有

  __global__ void kernel() {
     __shared__ int i;
     foo(i);
  }
  

  为什么不等价

  __global__ void kernel(int *i_ptr) {
     foo(*i_ptr);
  }
  
  int main() {
     int *i_ptr;
     cudaMalloc(&i_ptr, sizeof(int));
     kernel<<<blocks,threads>>>(i_ptr);
  }
  

关于全局内存与共享内存中的特定速度问题存在很多问题，但没有一个包含在实践中何时使用其中任何一个的概述。

非常感谢

Answer 1

• 当我们使用 cudaMalloc()

  为了在 gpu 上存储可以与主机通信的数据，我们需要分配的内存在它被释放之前一直存在，将全局内存视为在应用程序关闭或被释放之前一直存在的堆空间，它对任何具有指向该内存区域的指针的线程和 block 都是可见的。共享内存可以被视为具有生命的堆栈空间，直到内核的一个 block 完成，可见性仅限于同一 block 内的线程。所以cudaMalloc用于在全局内存中分配空间。

• 我们是否获得了指向共享内存或全局内存的指针？

  您将获得一个指向驻留在全局内存中的内存地址的指针。

• 全局内存是驻留在主机还是设备上？

  全局内存驻留在设备上。但是，有一些方法可以使用映射内存将主机内存用作“全局”内存，请参阅:CUDA Zero Copy memory considerations但是，由于总线传输速度的限制，它可能会很慢。

• 两者都有大小限制吗？

  全局内存的大小取决于卡与卡，从无到 32GB (V100)。而共享内存取决于计算能力。低于计算能力 2.x 的任何东西都具有每个多处理器最大 16KB 的共享内存(其中多处理器的数量因卡而异)。计算能力为 2.x 或更高的卡的每个多处理器至少有 48KB 的共享内存。

  见 https://en.wikipedia.org/wiki/CUDA#Version_features_and_specifications

  如果您使用的是映射内存，唯一的限制是主机有多少内存。

• 哪个访问速度更快？

  就原始数据而言，共享内存要快得多(共享内存~1.7TB/s，而全局内存~XXXGB/s)。然而，为了做任何你需要用一些东西填充共享内存的事情，你通常从全局内存中提取。如果对全局内存的内存访问是合并的(非随机的)和大字长，则根据卡及其内存接口(interface)，您可以实现接近数百 GB/s 的理论极限的速度。

  共享内存的使用是当您需要在一个线程 block 内重用已经从全局内存中提取或评估的数据时。因此，您无需再次从全局内存中提取，而是将其放入共享内存中，以供同一 block 内的其他线程查看和重用。

  也很常见用作便笺簿，以减少影响可以同时运行多少个工作组的寄存器压力。

• 在共享内存中存储变量是否与通过内核传递其地址相同？

  不，如果您传递任何东西的地址，它始终是全局内存的地址。从主机您无法设置共享内存，除非您将其作为常量传递，内核将共享内存设置为该常量，或者将地址传递给内核在需要时将其拉出的全局内存。