CUDA编程模型下，流式多处理器(SMP)活动块行为及声明验证

1. 分配给多处理器的活动块具有怎样的行为？ #

当一个活动块被分配到SM（流式多处理器）后，首先会被拆分成32线程一组的线程束（warp）来执行。SM会通过硬件调度器切换不同线程束的执行时机——这主要是为了隐藏内存访问的延迟，比如当某个线程束在等待内存数据时，SM就会调度其他线程束干活。

块内的所有线程共享该块专属的共享内存、寄存器资源，而且像__syncthreads()这种块级同步原语，只能作用于当前块内的所有线程，没法跨块同步。另外，一旦块被敲定分配到某个SM，它就会一直待在这个SM上直到执行完毕，期间绝不会被转移到其他SM上——毕竟块的状态和资源都绑定在当前SM上，迁移的成本太高了。

SM会维护自己的“活动块集合”，只有当某个活动块执行完、释放了占用的资源后，SM才会从全局的待分配块队列里捞新的块进来执行。

2. 针对SM执行的活动块行为及CUDA编程模型保障，各陈述判断与分析 #

我们逐个来看这几个陈述：

• 陈述1：当块被分配给SM并开始执行后，绝不会被转移至其他SM
  完全成立。CUDA编程模型明确保障了块的“绑定性”——一旦块被调度到某个SM，它的整个执行生命周期就和这个SM绑定了，直到所有线程完成任务。跨SM迁移块需要复制共享内存、寄存器状态等大量数据，开销大到不现实，CUDA模型也根本不需要这种机制来保证正确性。

• 陈述2：若核函数配置限制每个MP的活动集仅含N个块，则当SM活动集满N个块时，现有块完成前无法移除或添加新块
  成立。这里的“活动集”指的是SM上正在占用共享内存、寄存器等硬件资源的块。当活动集达到配置的N个上限时，SM已经没有多余资源来容纳新块了；而且正常执行情况下，SM也不会中途移除未完成的块（除非遇到错误终止）。只有等某个活动块执行完毕、释放了资源，SM才会尝试从全局队列里拿新的块加入活动集。

• 陈述3：若活动集中某块持续自旋，该SM上的其他块仍可继续执行
  成立。SM的调度是基于线程束的，不是针对整个块的。就算某个块里的线程束一直在自旋（比如死循环等操作），SM的硬件调度器会直接跳过它，去调度其他块的线程束执行。只要其他块已经在活动集里，它们的线程束就会正常被调度运行——不过要注意，如果自旋的线程束占用了太多寄存器资源，可能会限制SM能容纳的活动块总数，但不会影响已经在活动集里的其他块的执行。

• 陈述4：跨SM复制共享内存及执行状态、从全局内存备份/恢复并非合理方案，推测前两种行为受CUDA编程模型保障？
  这个推测是对的，前两种行为（块不跨SM迁移、活动集满时不添加/移除未完成块）确实是CUDA编程模型明确保障的。跨SM复制状态或备份到全局内存的方案之所以不合理，一是这种操作的性能开销极大，完全得不偿失；二是CUDA已经通过块绑定SM、活动集资源管理的机制，保证了块执行的正确性和效率，根本不需要这种复杂且低效的操作。

内容的提问来源于stack exchange，提问作者iam