并行编程实战——CUDA编程的Warp

最新推荐文章于 2025-11-16 10:27:24 发布

原创最新推荐文章于 2025-11-16 10:27:24 发布 · 505 阅读

10 ·

CC 4.0 BY-SA版权

文章标签：

#CUDA #c++

C++ 同时被 2 个专栏收录

455 篇文章

订阅专栏

并行编程

47 篇文章

订阅专栏

一、warp

warp，线程束。在前面的学习中，已经了解到CUDA中的执行调度单元是Warp而不是线程，一个Warp一般是有32个线程，一个Warp内的线程执行相同的指令。所以应该明白，上下文的切换是Warp而不是线程。这样的好处在于，可以更加合理的调配相关资源并发恢GPU的计算能力。这有点类似于，既考虑了整体的资源最优化又兼顾了同时多个线程应用的二者间一个平衡。
warp的调度是在CUDA中维护着一个WarpPool，可以依次从池中执行相关的Warp。所以SM的数量决定着Warp的最大并发度。同时，由于Warp的Context在内核启动时就已经存储在SM中，所以在调度时不需要进行上下文的的切换，只要选择合适的Warp即可。这是不是就是算法中的空间换时间的一个例子？当然，上下文的切换开销小不代表整体的开销小，这就需要开发者在开发时根据实际情况进行具体的分析处理。
当Warp成为Active Warp后，其可以开成三种情况：

select warp
选中的线程束将开始执行，并进行相关的任务操作
stalled warp
没有准备好的Warp，它将不会执行
eligible warp
已准备好但尚未执行的Warp

如果有OS中进程调度的知识的开发者可能会非常容易理解。

二、warp与内存

Warp级别的通信操作更轻量和方便，特别是相关内存的处理更是如此。对Warp中内存的使用可以通过下面的几个种情况来分析：

内存对齐
在同一个Warp中，在访问全局内存时，如果内存布局不合理，会导致时间的浪费，从而降低其效率。特别是当多个线程同时访问同一个bank时，会导致冲突。这时可以使用填充占用的方式进行内存的处理，避免冲突的可能。
通过设计使得Warp访问相同单元的内存
其实就是利用提高命中的机会和降低不同页之间操作的复杂度来提高效率和性能。

三、warp与块

Warp内线程的分配是受块的影响的，也就是说，块是分配到SM上的单元。如果块内线程的分配不合适，则会导致资源的不合理使用。比如有两个块，一个块是40个线程。那么这个块就会分配32个线程到一个Warp，而余下的8个线程分配到另外一个Warp中。这样，就会有两Warp产生了资源的浪费（余下的闲置的线程仍然会占用资源）。这和内存对齐的意思有些雷同。所以从原则上来讲，块线程的大小最好以Warp中线程容量的倍数来分配。
注意：虽然块可以有多个维度进行处理，但和数组类似，它其实是一个逻辑概念。在实际的物理操作中，仍然可以认为是一条线性的关系。

四、warp与线程及协作组

Warp一般是包含32个线程，这些线程组合在一起，从抽象的逻辑看，可以把它们看一个大的线程。它们步调一致，工作的任务相同（但数据不两只）。如果其中一个线程出现问题，则必然会影响到Warp整体的执行效率。
Warp内的数据交换可以使用warp shuffle，它可以直接使用寄存器的方式也可以通过Lane广播的方式到其它的Lane。在CUDA提供了原语层的支持：

__shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize)
__shfl_sync（）本质允许在不使用共享内存的情况下在warp内的线程之间交换变量。交换同时发生在warp中的所有活动线程中（并在掩码中命名），根据类型，每个线程移动4或8字节的数据。 Warp中的线程被称为lanes，它们的索引可能在0到warpSize-1（包括）之间。
它还有其它几种变种模式：__shfl_up_sync(mask, var, delta)，__shfl_down_sync(mask, var, delta)和__shfl_xor_sync(mask, var, laneMask)。具体的细节可参看CUDA的文档说明。
其它warp级别原语还有：
__ballot_sync（unsigned mask, predicate）
为掩码中所有未退出的线程计算谓词，并返回一个整数，当且仅当第n个线程的谓词计算结果为非零且第n个线程处于活动状态时，该整数的第n位被设置
__any_sync(unsigned mask, predicate) / __all_sync(unsigned mask, predicate)
warp中至少有一个(any)或所有(all)线程的谓词（predicate）为true返回非0，否则返回0。
__activemask
获取当前Warp中的有效线程的掩码，用于和__shfl_sync等一起进行Warp的相关操作

协作组可以在warp中平铺线程以支持warp级别的协作操作。即如果tile与warp的大小匹配，CUDA则可以省略warps的隐式同步，确保正确的内存操作以避免竞争条件。通过这种方式，消除了隐式同步，从而增强GPU的性能。

五、Warp Divergence

线程否的分歧或线程束的发散，其实就是控制分支。在开发过程中，一定会出现不同（控制流语句）逻辑结果问题，这样分支就出现了。
首先要明白latency的内容，简单说就是指令从开始到真正开始执行的时钟周期数。学硬件都知道，高电平和低电平的的出现是一个过程，而不是在一个点完成的，它有一个时间过渡过程。所以在硬件的文档（Datasheet）中，高低电平的判断需要一个延时。大家可以把latency理解成这个延时过程。它分为两种情况：

Arithmetic instruction
算术指令操作的始末时间间隔
Memory instruction
内存指令操作的时间间隔

在同一个Warp中，线程执行相同的指令。对于Warp中的线程来说，处理分支时，如果有线程计算结果为True，则并行执行为True的线程，计算结果为False的线程等待（stalled）前者执行完成后，再执行。这等于是浪费另外一个逻辑层次的所有线程。
为了提高利用率，分支预测就引入了。分支预测技术已经在各个方面被广泛应用，此处就不再赘述，如果不明白可以查看相关的技术资料。
另外，为了进一步提高利用率，CUDA引入延迟隐藏，这个技术其实非常好理解，大家在学习CPU的指令流水时的多个指令流的操作和这个非常类似。刚刚说过，一个Warp的Context都是相同的，而每个SM中的Warp中Context都是片上的，即其切换的代价可以忽略。那么就可以将eligible warp填充到执行的Wrap的latency指令中即通过指令填充形成一条连续流水，这样可以最大化的降低指令执行周期的等待。