如果SM上一个warp被stalled，warp scheduler会调度一个candidate warp来替换被stalled的warp以掩藏latency。那么，这个candidate warp必须来自和stalled warp同一个block呢，还是必须来自另一个不同的block，还是二者同等地位，两种情况都可以？
这个应该没有限制，只要是当前加载于该SM上的block里面就绪的warp应该都可以。
我同意2楼版主的说法, 应该是SM上的warp只要就绪就能够被选入执行
我最开始也是这样想，后来前几天老师上课讲掩藏latency的方法中有一种是适当减小threadblock的规模，为什么要这样做呢？当时忘记老师的解释了:L，于是就猜，如果计算规模和总的线程数目不变，减小threadblock会使block数目增多，但是warp数也没有增多，这对掩藏latency有什么好处呢？除非能入选执行的warp来自不同的block，这样的话，block数目增多，来自不同block的warp数目也增多，就可以掩藏延迟了。
但是楼上两位的解释又让我把这个想法否定了，因为如果没有减小threadblock规模这种优化方法，这样做本来就没有道理。
现在又想了想，是不是跟__syncthreads()有关，如果kernel中有同步操作的话，那么相同block中的warp难免会受到影响而停在同步点无法掩藏延迟：
memory
__syncthreads(); //barrier
math
math
…
stall
但不同block中的warp就不受这个限制了，可以不用等其他block的barrier而一直执行下去，做到掩藏延迟。所以在有同步操作的kernel中，适当减小threadblock规模以增加block数目才有意义。不知是不是这样？
我觉得可以考虑一些特殊情况来辅助判断。
比如说，你一个block有1024个线程，这样一个本来能容纳1536个线程的SM就只能容纳一个block，也就是只放置了1024个线程，线程数量减少对于掩盖延迟是不利的。
或者还是上面的例子，如果这1024个线程同属一个block，此时block内有syncthreads，若线程运行到这里，那必须停下来等待后面的线程；而此时后面的线程如果遇到长延迟的操作，将不再有线程可以切换而只能等待。而此时如果拆成2个512线程的block，各维护各的syncthreads，两边一般不是同步的，可以互相掩盖。
欢迎莅临cudazone，祝您编码愉快。
对阿 确实如此 您的解释我很认同。
受教了！
欢迎您莅临cudazone积极讨论问题，祝您编码愉快~