虽然低占用率的SM确实无法充分隐藏延迟,但重要的是要理解这一点:
更高的占用率!=更高的吞吐量!
占用率只是衡量SM在任何特定时刻可供选择的工作量的指标.具有更多驻留扭曲使SM更有能力执行有用的工作,而其他warp正在等待结果(内存访问的结果或计算 - 都具有非零延迟).
吞吐量是每秒完成工作量的量度,虽然它可能受到延迟(因此占用率)的限制,但它也可能受到内存带宽,指令吞吐量(执行单元数量)和其他因素的限制.
编程指南说明拥有多个线程块而不仅仅是一个大线程块更好的原因是因为有时候能够不仅从其他warp而且从其他块发出工作更好.这是一个例子:
想象一下,你的大线程块必须从全局内存加载数据(高延迟)并将其存储到共享内存(低延迟),然后必须立即执行__syncthreads().在这种情况下,当warp完成加载其数据并将其写入共享内存时,它必须等待,直到块中的所有其他线程完成相同操作.对于大块,可能需要一段时间.但是如果有多个较小的线程块占用SM,那么SM可以在等待__syncthreads第一个块中的满足时从其他块切换并工作.这有助于减少GPU空闲时间并提高效率.
您不一定非常希望拥有非常小的块(因为Fermi上的SM支持最多8个驻留块),但是拥有128-512个线程的块通常比使用1024个线程的块更有效.