为什么还要了解CUDA Warps？

Question

我有GeForce GTX460 SE,所以它是:6 SM x 48 CUDA核心= 288 CUDA核心.众所周知,在一个Warp中包含32个线程,并且在一个块中同时(一次)只能执行一个Warp.也就是说,在单个多处理器(SM)中,即使有48个可用核心,也可以同时只执行一个Block,一个Warp和32个线程？

此外,可以使用threadIdx.x和blockIdx.x来分发具体的Thread和Block的示例.要分配它们,请使用内核<<< Blocks,Threads >>>().但是如何分配特定数量的Warp-s并分发它们,如果不可能那么为什么还要去了解Warps呢？

Answer 1

这种情况比你描述的要复杂得多.

ALU(核心),加载/存储(LD/ST)单元和特殊功能单元(SFU)(图像中为绿色)是流水线单元.它们在完成的各个阶段同时保留许多计算或操作的结果.因此,在一个周期内,他们可以接受新操作并提供很久以前开始的另一个操作的结果(如果我没记错的话,ALU大约需要20个周期).因此,理论上单个SM具有用于同时处理48*20个周期= 960个ALU操作的资源,即每个warp = 30个warp的960/32个线程.此外,它可以处理LD/ST操作和SFU操作,无论其延迟和吞吐量如何.

warp调度程序(图像中的黄色)可以为每个warp = 64个线程调度2*32个线程到每个循环的管道.这就是每个时钟可以获得的结果数量.因此,考虑到混合的计算资源,48核心,16 LD/ST,8 SFU,每个具有不同的延迟,同时处理混合的warp.在任何给定的周期,warp调度程序尝试"配对"两个warp以进行调度,以最大化SM的利用率.

如果指令是独立的,则warp调度程序可以从不同的块或同一块中的不同位置发出warp.因此,可以同时处理来自多个块的warp.

增加复杂性,执行少于32个资源的指令的warp必须多次发出才能为所有线程提供服务.例如,有8个SFU,因此这意味着包含需要SFU的指令的warp必须被调度4次.

该描述被简化.还有其他限制可以确定GPU如何安排工作.您可以在网上搜索"fermi architecture"来查找更多信息.

所以,来到你的实际问题,

为什么懒得去了解Warps？

当您尝试最大化算法的性能时,了解warp中的线程数并将其考虑在内变得非常重要.如果您不遵守这些规则,则会失去性能:

• 在内核调用中<<<Blocks, Threads>>>,尝试选择一些线程,这些线程与warp中的线程数均分.如果不这样做,最终会启动包含非活动线程的块.

• 在您的内核中,尝试让warp中的每个线程遵循相同的代码路径.如果不这样做,就会得到所谓的经线发散.这是因为GPU必须通过每个不同的代码路径运行整个warp.

• 在您的内核中,尝试将每个线程置于warp加载中并以特定模式存储数据.例如,让warp中的线程在全局内存中访问连续的32位字.