42

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

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

  • 将变量存储在共享内存中与通过内核传递其地址相同吗？即而不是拥有

    __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);
}

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

非常感谢

4

3 回答 3

56

  • 当我们使用 cudaMalloc()

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

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

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

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

    全局内存驻留在设备上。但是，有一些方法可以使用映射内存将主机内存用作“全局”内存，请参阅：CUDA 零拷贝内存注意事项，但是，由于总线传输速度的限制，它可能会降低速度。

  • 任何一个都有大小限制吗？

    全局内存的大小取决于卡与卡，从无到 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 的理论极限的速度，具体取决于卡及其内存接口。

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

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

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

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

于 2012-12-30T19:28:38.787 回答
11

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

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

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

更新：

Compute Capability 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>>>(...);
于 2012-12-30T19:25:11.107 回答
3

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

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

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

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

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

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

于 2012-12-30T19:19:45.770 回答