cuda - 基本 cuda 共享内存

Question

我是 cuda 新手，对共享内存有一些疑问：

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

2. SM 如何在块之间划分共享内存？它是平均分配的（例如，如果有 2 个块，那么每个块在 SM 中获得一半的共享内存，而不管实际使用了多少），还是基于需要？

3. 我对共享内存bank的理解是：共享内存分为32个同样大的内存bank。那么这意味着每个区块（即每个区块都有自己的 32 个银行）还是每个 SM？

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

谢谢！

标签： cuda

Answer 1

首先让我指出，共享内存首先是编程模型的抽象，通过它可以暴露硬件的某些特性（快速的片上内存）。在 CUDA 编程模型中，网格中的每个块（内核启动）都获得相同数量的共享内存。多少取决于内核函数所需的静态分配共享内存量以及内核启动中指定的任何额外动态共享内存。

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

是的，目前就是这样。但是，这与您对 CUDA 的编程方式并不像您想象的那样真正相关，因为：

2. SM 如何在块之间划分共享内存？它是平均分配的（例如，如果有 2 个块，那么每个块在 SM 中获得一半的共享内存，而不管实际使用了多少），还是基于需要？

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

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

3. 我对共享内存bank的理解是：共享内存分为32个同样大的内存bank。那么这意味着每个区块（即每个区块都有自己的 32 个银行）还是每个 SM？

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

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

你不能cudaMemcpy()进入共享内存。共享内存只能由同一块的设备线程访问，并且只要该块正在运行，它就会持续存在。