CUDA共享内存和全局内存有什么区别?

46

我对如何在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);
    }
    

关于全局内存和共享内存的特定速度问题,已经有很多问题,但没有一个涵盖实践中何时使用其中任何一个的概述。

非常感谢。


1
@NolwennLeGuen 问题在于你无法控制L1-L2缓存中驻留的数据。而在共享内存中,你清楚地知道那里面有什么。 - 1-----1
3个回答

61
  • 当我们使用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的理论极限速度,具体取决于卡和其内存接口。

    使用共享内存的情况是在块内线程中,重复使用已从全局内存中提取或评估的数据。因此,您可以将其放入共享内存中供同一块内的其他线程查看和重用,而不是再次从全局内存中提取。

    它也经常用作临时缓冲区,以减少寄存器压力,影响可以同时运行的工作组数量。

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

    不是。如果您传递任何东西的地址,它总是指向全局内存的地址。从主机上,您无法设置共享内存,除非将其作为常量传递给内核,在这种情况下,内核将共享内存设置为该常量,或者将其作为指向全局内存的地址传递,内核需要时会从中提取。


12
全局内存的内容对网格中的所有线程可见。任何线程都可以读写全局内存的任何位置。
共享内存是网格中每个块独立的。任何一个块的线程都可以读写该块的共享内存。一个块中的线程不能访问另一个块的共享内存。
1. cudaMalloc始终分配全局内存。 2. 全局内存在设备上。 3. 显然,每个内存都有一个大小限制。全局内存是您使用的GPU的DRAM的总量。例如,我使用的GTX460M具有1536 MB DRAM,因此拥有1536 MB的全局内存。共享内存由设备架构指定,并以每个块为基础进行测量。计算能力为1.0至1.3的设备具有16 KB/Block,从2.0开始的计算具有48 KB/Block的默认共享内存。 4. 共享内存访问速度比全局内存快得多。它就像在块的线程之间共享的本地缓存。 5. 不,只有全局内存地址可以传递到从主机启动的内核。在你的第一个示例中,变量是从共享内存中读取的,而在第二个示例中,它是从全局内存中读取的。
更新:
计算能力为7.0的设备(Volta架构)允许每个块分配最多96 KB的共享内存,前提是满足以下条件。
1. 动态地分配共享内存。 2. 在启动内核之前,使用函数cudaFuncSetAttribute指定动态共享内存的最大大小,如下所示。
__global__ void MyKernel(...)
{
    extern __shared__ float shMem[];
}

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

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

3
CUDA共享内存是该块中线程之间共享的内存,即在网格中的块之间,共享内存的内容是未定义的。可以将其视为手动管理的L2缓存。
通常情况下,全局内存位于设备上,但如果设备支持,则最近版本的CUDA可以将主机内存映射到设备地址空间,从而在这种情况下触发主机到设备内存的原地DMA传输。
共享内存存在大小限制,取决于设备。它在枚举CUDA设备时检索的设备功能中报告。全局内存受GPU可用总内存的限制。例如,GTX680提供48kiB的共享内存和2GiB的设备内存。
与全局内存相比,共享内存访问速度更快,但是必须仔细对齐访问模式(对于共享和全局内存)才能有效。如果无法使访问模式正确对齐,请使用纹理(还是全局内存,但通过不同的电路和缓存访问,可以更好地处理不对齐的访问)。

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

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

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接