2

我阅读了一些引用本地内存的 CUDA 文档。（主要是早期文档。）设备属性报告本地内存大小（每个线程）。“本地”内存是什么意思？什么是“本地”内存？“本地”内存在哪里？如何访问“本地”内存？是__device__记忆，不是吗？

设备属性还报告：全局、共享和恒定的内存大小。这些陈述是否正确： 全局内存就是__device__内存。它具有网格范围和网格（内核）的生命周期。 不变的记忆就是__device__ __constant__记忆。它具有网格范围和网格（内核）的生命周期。 共享内存就是__device__ __shared__内存。它具有单个块范围和该块（线程）的生命周期。

我认为共享内存是 SM 内存。即只有那个单一的SM可以直接访问的内存。相当有限的资源。SM不是一次分配一堆块吗？这是否意味着 SM 可以交错执行不同的块（或不能）？即运行块* A * 线程，直到它们停止。然后运行 ​​block* B * 线程直到它们停止。然后再换回 block* A * 线程。或 SM 是否为 block* A * 运行一组线程，直到它们停止。然后换入另一组块* A * 线程。此交换继续进行，直到块* A * 用尽。然后才开始在块* B上工作*。我问是因为共享内存。如果单个 SM 正在从 2 个不同的块交换代码，那么 SM 如何快速换入/换出共享内存块？（我认为后面的 senerio 是真的，并且没有换入/换出共享内存空间。块 * A * 运行直到完成，然后块 * B * 开始执行。注意：块 * A * 可能是不同的内核比块* B *。）

4

1 回答 1

5

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

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

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

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

于 2012-08-02T20:30:40.313 回答