cuda 共享内存和块执行调度

Posted 2023-04-15

【中文标题】cuda 共享内存和块执行调度

【英文标题】：cuda shared memory and block execution scheduling

【发布时间】：2012-09-21 01:01:37

【问题描述】：

我想根据每个块使用的共享内存量来清除 CUDA 共享内存 和 块执行 的执行状态。

状态

我的目标是 GTX480 nvidia 卡，每块有 48KB 共享内存和 15 个流式多处理器。所以，如果我声明一个包含 15 个块的内核，每个块使用 48KB 的共享内存，并且没有达到其他限制（寄存器、每个块的最大线程数等），每个块都运行到一个 SM（15 个）直到结束。在这种情况下，只需要在同一块的 warp 之间进行调度。

问题

所以，我的误解是： 我调用具有 30 个块的内核，以便每个 SM 上驻留 2 个块。现在每个 SM 上的 scheduler 必须处理来自不同块的扭曲。但只有当一个块完成执行时，另一个块的扭曲才会在 SM 上执行，因为共享内存总量（每个 SM 48KB）使用。如果这没有发生并且不同块调度在同​​一个 SM 上执行的扭曲，结果可能是错误的，因为一个块可以读取从另一个块加载到共享内存中的值。我说的对吗？

【参考方案1】：

您无需担心这一点。正如您所说的那样，如果由于使用的共享内存量而每个 SM 只适合一个块，那么任何时候都只会安排一个块。因此，不会因为过度使用共享内存而导致内存损坏。

---

顺便说一句，出于性能原因，通常每个 SM 至少运行两个块会更好，因为

在 __syncthreads() 期间，SM 可能会不必要地空闲，因为来自块的越来越少的 warp 仍然可以运行。 同一块的warp 倾向于紧密耦合运行，因此有时所有warp 都在等待内存，而其他时候所有warp 都执行计算。如果使用更多块，这可能会更好，从而提高整体资源利用率。

当然，与每个 SM 运行多个块相比，每个块的共享内存越多可以提供更大的加速，这可能是有原因的。