1

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

状态

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

问题

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

4

1 回答 1

2

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

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

  • 在 __syncthreads() 期间，SM 可能会不必要地空闲，因为来自块的越来越少的扭曲可能仍然可以运行。
  • 同一块的 warp 往往是紧密耦合的，所以有时所有的 warp 都在等待内存，而其他时候所有的 warp 都执行计算。有了更多的块，这可能会更好，从而整体上更好地利用资源。

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

于 2012-09-29T11:04:52.867 回答