本地、全局、常量和共享内存

Question

我读了一些CUDA文档，其中提到了本地内存。（这主要是早期的文档。）设备属性报告了一个本地内存大小（每个线程）。'本地'内存是什么意思？'本地'内存是什么？'本地'内存在哪里？如何访问'本地'内存？它是__device__内存，是吗？

设备属性还报告了全局、共享和常量内存大小。这些说法正确吗： 全局内存是__device__内存。它具有网格范围，并且在网格（内核）的生命周期内存在。 常量内存是__device__ __constant__内存。它具有网格范围和网格（内核）的生命周期。 共享内存是__device__ __shared__内存。它具有单个块范围，生命周期为该块（线程）。

我认为共享内存是SM内存。即只有该单个SM可以直接访问的内存。这是一种相当受限制的资源。一个SM同时被分配了一堆块吗？这是否意味着SM可以交错执行不同的块（或不执行）？即运行块*A*线程直到它们停止。然后运行块*B*线程，直到它们停止。然后再切换回块*A*线程。或者，SM是否运行一组线程（对于块*A*），直到它们停止。然后另一组块*A*线程被交换进来。这种交换会继续进行，直到块*A*耗尽。只有在此时才开始在块*B*上工作。 我之所以问这个问题是因为共享内存。如果单个SM正在从2个不同的块中交换代码，那么该SM如何快速地交换共享内存块？（我认为后一种情况是正确的，并且没有在共享内存空间中交换。块*A*运行直到完成，然后块*B*开始执行。注意：块*A*可能是不同于块*B*的内核。）

- Doug

---

“local”在“本地内存”中指的是“线程本地”，即每个线程都操作自己的实例。你说得对，这只是从总体__device__内存中划分出来的内存。相比之下，__shared__内存对同一线程块中的所有线程都是可访问的，即它是由线程块中的线程共享的。全局内存可以被整个网格中的任何线程访问，但是如果不同的线程想要访问相同的位置，必须小心处理，例如通过使用原子操作。 - njuffa

Answer 1

从CUDA C编程指南5.3.2.2节中，我们可以看到本地内存在以下几种情况下使用：

• 当每个线程有一些数组但它们的大小在编译时未知（因此可能无法适应寄存器）
• 当数组的大小在编译时已知，并且该大小对于寄存器内存来说太大（这也可能发生在大型结构体中）
• 当内核已经使用完所有寄存器内存时（因此如果我们用n int填满了寄存器，则第n+1个int将进入本地内存）-这最后一种情况是寄存器溢出，应该避免，因为：

“本地”内存实际上位于全局内存空间中，这意味着与之进行读写相比，寄存器和共享内存的速度要慢得多。每次在内核中使用某个变量、数组等不适合寄存器、不是共享内存并且没有作为全局内存传递的东西时，都会访问本地内存。你不必做任何明确的事情来使用它-事实上，你应该尽量减少它的使用，因为寄存器和共享内存要快得多。

编辑： 关于共享内存，您不能让两个块交换共享内存或查看彼此的共享内存。由于块的执行顺序不能保证，如果您尝试这样做，您可能会占用SMP数小时等待另一个块被执行。同样，正在设备上运行的两个内核不能看到彼此的内存，除非它是全局内存，即使是这样，您也在冒竞争条件的风险。据我所知，块/内核实际上无法相互发送“消息”。您的情况并没有真正意义，因为块的执行顺序每次都会不同，而且等待另一个块是不好的做法。

- limes