cuda - 基本cuda共享内存

Question

我是 CUDA 新手，有一些关于共享内存的问题:

1. 每个 SM 在同一 GPU 中是否具有相同数量的共享内存？

2. SM 如何在 block 之间划分共享内存？是否平均分配(例如，如果有 2 个 block ，则无论实际使用多少，每个 block 都会获得 SM 内一半的共享内存)，还是根据需要分配？

3. 我对共享内存Bank的理解是:共享内存分为32个同样大的内存Bank。那么这意味着每个区 block (即每个区 block 都有自己的 32 个银行)还是每个 SM？

4. 如果我在多个字的共享内存中执行 cudaMemcpy，这算作单个事务还是多个事务？这会导致银行冲突吗？

谢谢!

Answer 1

首先让我指出，共享内存首先是编程模型的抽象，通过该模型可以暴露硬件的某些功能(快速的片上内存)。在 CUDA 编程模型中，网格(内核启动)中的每个 block 都获得相同数量的共享内存。具体多少取决于内核函数所需的静态分配共享内存的数量以及内核启动中指定的任何其他动态共享内存。

1. does every SM have the same amount of shared memory within the same GPU?

是的，目前就是这种情况。然而，这与您对 CUDA 编程的方式并不像您想象的那么相关，因为:

2. How does an SM partition the shared memory amongst the blocks? Is it distributed equally ( ex. if there are 2 blocks, then each block gets half the shared memory within the SM regardless of how much is actually used ), or is it based on the needs?

当您启动内核时，您可以指定每个 block 需要多少共享内存。然后，这会告知每个多处理器上可以容纳多少个 block 。因此，并不是 block 的数量定义了每个 block 获得多少共享内存，而是相反:每个 block 所需的共享内存量是定义每个多处理器上可以驻留多少 block 的因素之一。

您将需要阅读有关延迟隐藏和占用的内容，因为在 GPU 编程方面这些是非常基本的主题。有关不同 GPU 架构的内存子系统的更多详细信息，请查看 CUDA Programming Guide .

3. My understanding of a shared memory bank is: shared memory is divided into 32 equally large memory banks. So does this mean per block ( i.e. eveyr block has their own 32 banks ) or is it per SM?

最后，由于 GPU 核心的 SIMD (SIMT) 性质，实际的程序执行发生在 warp 中。当这样的 warp(目前，这实际上意味着一组 32 个线程)执行共享内存访问时，当服务该指令生成的共享内存请求时，存储体冲突将成为一个问题。是否可以并行处理多个 warp 的共享内存请求并没有真正记录。我的猜测是，每个 SM 只有一个单元来处理共享内存请求，因此答案是否定的。

4. If I perform a CudaMemcpy from / into shared memroy of more than one word, does this count as a single transaction or multiple transactions? And could this cause bank conflicts?

您无法cudaMemcpy()进入共享内存。共享内存只能由同一 block 的设备线程访问，并且仅在该 block 运行时持续存在。