1、CUDA_C_BEST_PRACTIES_GUIDE中讲到，要隐藏一条warp的24周期延迟，在1.X上就需要248 = 192个active threads。而在2.0上就需要2432 = 768个active threads。我不知道这些延迟具体是怎么被隐藏起来的，麻烦大牛们能详细说明一下:)。
2、书上有讲在1.0/1.1设备上，一个SM可以保存24个warp（768threads），1.2/1.3则可以保存32个warp（1024threads），在2.0上则更高。我想问一下，一个SM可以保存的warp总量具体是根据什么得来的？有没有什么具体的计算方式，还是说根据硬件能力而固定的:)？
3、CUDA_C_BEST_PRACTIES_GUIDE和书上有讲到，提高效率，每个SM至少得拥有2个block，这是根据什么得来的？仅仅是为了减少block中线程间同步（__syncthreads）的时间？
谢谢:)
- 当一个warp出现延迟时，如果有其它的warp中的线程处于active状态的话，这些active的线程可以执行，这样可以隐藏延迟，就是充分利用有限的资源，抓住一分一秒。
2.一个SM保存的warp数量是由一个SM加载的最大线程数量计算的，这与硬件的计算能力有关。
3.和1有点相似，至少两个block的话，当一个block出现延迟的时候，可以由另外一个block使用硬件资源
谢谢你的回答:)
其实关于你回答的第一点，我也知道。我只是想清楚的了解其隐藏延迟的流程。比如说现只有两个active warp，第一个warp出现延迟，那第二个warp会跟上执行，占用计算资源。那当第一个warp度过延迟后，是等待第二个warp出现延迟后它补上去执行，还是说不管第二个warp有没有在执行，就立马把计算资源让出来呢？
第二点中，我还是不清楚在1.X和2.X设备上怎么计算SM加载的最大线程数，还得再麻烦你一次。
1.是第一种情况即当第一个warp等待结束并且另外一个warp出现等待或者计算完毕之后才会继续执行。
2.理论最大线程数是由硬件决定的，比如1.x是768,2.x是1536，但是实际加载的线程数不超理论值，这主要是受到SM中寄存器，共享存储器，block数的限制，这些限制你可以在书上看到的。比如一个block你使用的shared达到了SM中的最大值，那么最多就加载一个block，里面可能仅有256个线程。
谢谢你的回答:)。关于第二个问题，是不是可以理解为能装载的理论最大线程数是随着寄存器和shared的增大而增大?
这个是由硬件决定的，也不能那么说，因为寄存器或者shared增大之后，可以意味着每个线程使用的最大资源有所提高，所以不一定意味着线程理论值增大啊。呵呵，所以，一个SM最大加载的线程是由硬件的计算能力决定的，不必纠结。只要知道各种资源的限制，在使用中都尽量不超过其上限即可。
:)非常感谢…
机制：当一个warp在等待的时候，调度器就去调度另一个不用等待的warp来执行。
一个SM的warp上限是硬件决定的。
主要是考虑__syncthreads()的开销。
看了你们的讨论，受益了，谢谢:)