cuda共享内存和块执行调度



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

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

问题

因此,我的误解是:
我调用一个有30个块的内核,因此每个SM上有2个块。现在,每个SM上的调度器必须处理来自不同块的扭曲。但只有当一个块完成执行时,由于共享内存总量(每个SM 48KB)的使用,另一个块的扭曲才会在SM上执行。如果这种情况没有发生,并且调度在同一SM上执行的不同块发生扭曲,则结果可能是错误的,因为一个块可以读取共享内存中从另一个块加载的值。我说得对吗?

您不需要担心这一点。正如您正确地说的,如果由于使用的共享内存量,每个SM只适合一个块,那么在任何时候都只能调度一个块。因此,不存在由于过度使用共享内存而导致内存损坏的可能性。


BTW出于性能原因,通常每个SM至少运行两个块更好,因为

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

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