CUDA 共享内存和全局内存有什么区别？答案

【问题标题】：What's the difference between CUDA shared and global memory?CUDA 共享内存和全局内存有什么区别？
【发布时间】：2012-12-15 03:30:37
【问题描述】：

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

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

正在存储一个共享内存中的变量与通过核心？ 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);
}

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

非常感谢

【问题讨论】：

@NolwennLeGuen 问题是您无法控制 L1-L2 缓存中的数据。在共享内存中，您确切知道那里有什么。

标签： memory cuda global shared-memory

【解决方案1】：

全局内存的内容对网格的所有线程都是可见的。任何线程都可以读写全局内存的任何位置。

共享内存对于网格的每个块都是独立的。块的任何线程都可以读取和写入该块的共享内存。一个块中的线程不能访问另一个块的共享内存。

cudaMalloc 总是分配全局内存。
全局内存驻留在设备上。
显然，每个内存都有大小限制。全局内存是您正在使用的 GPU 的 DRAM 总量。例如，我使用具有 1536 MB DRAM 的 GTX460M，因此具有 1536 MB 全局内存。共享内存由设备架构指定，并基于每个块进行测量。计算能力 1.0 到 1.3 的设备默认具有 16 KB/Block，计算 2.0 及更高版本默认具有 48 KB/Block 共享内存。
共享内存的访问速度比全局内存快很多。它就像一个块的线程之间共享的本地缓存。
没有。只有全局内存地址可以传递给从主机启动的内核。在您的第一个示例中，变量是从共享内存中读取的，而在第二个示例中，它是从全局内存中读取的。

更新：

如果满足以下条件，计算能力 7.0（Volta 架构）的设备允许为每个块分配高达 96 KB 的共享内存。

共享内存是动态分配的
在启动内核之前，使用函数cudaFuncSetAttribute 指定动态共享内存的最大大小，如下所示。

__global__ void MyKernel(...)
{
    extern __shared__ float shMem[];
}

int bytes = 98304; //96 KB
cudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes);

MyKernel<<<gridSize, blockSize, bytes>>>(...);

【讨论】：

【解决方案2】：

当我们使用 cudaMalloc() 时

为了在 gpu 上存储可以传回主机的数据，我们需要分配的内存在它被释放之前一直存在，将全局内存视为在应用程序关闭或被释放之前一直存在的堆空间，它对任何具有指向该内存区域的指针的线程和块都是可见的。共享内存可以被视为具有生命的堆栈空间，直到内核的一个块完成，可见性仅限于同一块内的线程。所以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）。然而，为了做任何你需要用一些东西填充共享内存的事情，你通常会从全局内存中拉取。如果对全局内存的内存访问是合并的（非随机的）和大字长，您可以实现接近数百 GB/s 的理论极限的速度，具体取决于卡及其内存接口。

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

也很常见用作便笺簿，以减少影响可以同时运行多少个工作组的寄存器压力。
在共享内存中存储变量是否与通过内核传递其地址相同？

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

【讨论】：

【解决方案3】：

CUDA 共享内存是块内线程之间共享的内存，即网格中的块之间共享内存的内容是未定义的。可以认为是手动管理的二级缓存。

通常全局内存驻留在设备上，但最新版本的 CUDA（如果设备支持）可以将主机内存映射到设备地址空间，在这种情况下触发从主机到设备内存的原位 DMA 传输。

共享内存有大小限制，具体取决于设备。它在设备功能中报告，在枚举 CUDA 设备时检索。全局内存受限于 GPU 可用的总内存。例如，GTX680 提供 48kiB 共享内存和 2GiB 设备内存。

共享内存的访问速度比全局内存快，但访问模式必须仔细对齐（对于共享内存和全局内存）以提高效率。如果您无法使访问模式正确对齐，请使用纹理（也是全局内存，但通过不同的循环和缓存访问，可以更好地处理未对齐的访问）。

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

不，绝对不是。您提出的代码将是您使用原位传输的全局内存的情况。共享内存不能在内核之间传递，因为共享块的内容仅在线程的执行块中定义。

【讨论】：