我是 cuda 新手，有几个关于共享内存的问题：
每个 SM 在同一个 GPU 中是否拥有相同数量的共享内存？
SM 如何在块之间划分共享内存？是平均分配（例如，如果有 2 个块，那么无论实际使用多少，每个块都获得 SM 内一半的共享内存），还是根据需要？
我对共享内存bank的理解是：共享内存分为32个同样大的内存bank。那么这意味着每个区块（即每个区块都有自己的 32 个银行）还是每个 SM？
如果我从/到多个单词的共享内存中执行 CudaMemcpy，这算作单个事务还是多个事务？这会导致银行冲突吗？
谢谢！
【问题讨论】：
标签： cuda
首先我要指出，共享内存首先是编程模型的抽象，通过它可以暴露硬件的某些特性（快速、片上内存）。在 CUDA 编程模型中，网格中的每个块（内核启动）都获得相同数量的共享内存。多少取决于内核函数所需的静态分配共享内存量以及内核启动时指定的任何额外动态共享内存量。
- 每个 SM 在同一个 GPU 中是否具有相同数量的共享内存？
是的，目前就是这样。但是，这与您对 CUDA 的编程方式并不像您想象的那样真正相关，因为：
- SM 如何在块之间划分共享内存？是平均分配（例如，如果有 2 个块，那么无论实际使用多少，每个块都获得 SM 内一半的共享内存），还是根据需要？
当您启动内核时，您指定每个块需要多少共享内存。然后，这会通知每个多处理器上可以容纳多少块。因此，并不是块的数量决定了每个块获得多少共享内存，而是反过来：每个块所需的共享内存量是定义每个多处理器上可以驻留多少块的因素之一。
您需要了解延迟隐藏和占用，因为这些是 GPU 编程的基本主题。有关不同 GPU 架构的内存子系统的更多详细信息，请查看CUDA Programming Guide。
- 我对共享内存bank的理解是：共享内存分为32个同样大的内存bank。那么这意味着每个区块（即每个区块都有自己的 32 个银行）还是每个 SM？
最后，由于 GPU 内核的 SIMD (SIMT) 特性，实际的程序执行发生在 warp 中。当这样的 warp（目前，这实际上意味着一组 32 个线程）执行共享内存访问时，由于该指令生成的共享内存请求得到服务，银行冲突将成为一个问题。是否可以并行处理多个 warp 的共享内存请求并没有真正记录。我的猜测是每个 SM 只有一个单元来处理共享内存请求，因此，答案是否定的。
- 如果我从/到多个单词的共享内存中执行 CudaMemcpy，这算作单个事务还是多个事务？这会导致银行冲突吗？
您不能cudaMemcpy() 进入共享内存。共享内存只能由同一块的设备线程访问，并且只在该块运行时持续存在。
【讨论】：