8.6 其他指令

8.6.1 线程束级原语

对于CUDA程序员来说，用不了多少时间就会意识到线程束（线程与线程块之间的一个基本单位）作为基本执行单元的重要性。在流处理器簇版本还是1.x的时代，英伟达公司就开始增加了一些仅作用于线程束的特殊指令。

1. 投票

CUDA架构是32位的，所以每一个线程束包含32个线程，是一个使指令能够评估条件并且广播1位结果给线程束中每一个线程的最适合的选择。VOTE指令（在SM 1.2中首次应用）就是用于评估条件并将1位结果广播给线程束中所有线程的。_any()函数在线程束32个线程中有任何一个判断为真时会返回1。_all()函数在线程束32个线程都判断为真时会返回1。

费米架构中增添了一个VOTE的变种，它可以传回一个线程束中32个线程的全部断定结果。函数__ballot()会为线程束中所有线程评估条件，并且返回一个32位的值，每一位分别代表着相同序号的线程的环境。

2. 洗牌

开普勒架构增添了洗牌（shuffle）指令，这个指令允许同一个线程束中的线程间进行数据交换而不经过共享内存中转。尽管这样的传输并没有节省时间，但这种方式有助于避免过多的读写操作，并且减少共享内存的使用。

下面的指令被包装进许多设备函数，使用了定义于sm_30_instrinsic.h中的内联PTX汇编代码。

int shfl(int var, int srcLane, int width=32);  
int shfl_up(int var, unsigned int delta, int width=32);  
int shfl_down(int var, unsigned int delta, int width=32);  
int shfl_xor(int var, int laneMask, int width=32);

宽度参数width，默认为线程束的宽度32，该值一定在 $2\sim 32$ 的范围内，并为2的幂次。它使线程束细分成段。如果width小于32，线程束的每段会作为一个单独的实体，并且起始逻辑线程编号为0。一个线程仅能与它在同一段的线程进行数据交换。

__shfl()返回的var值由编号为srcLane的线程给出。如果srcLane超过0到width-1的范围，线程本身的值var被返回。该指令变种可用来在一个线程束内广播值。__shfl_up()通过从调用者线程编号减去delta，并夹取到0到width-1的范围内，计算源线程编号。

_shfl_down()则通过在调用者线程编号上加上delta来计算源线程编号。

__shfl_up()和__shfl_down()分别启用了线程束级扫描和反转扫描操作。__shfl_xor()通过执行调用者线程编号与laneMask的按位异或计算源线程编号，存在源线程中的var值被返回。该指令变种可以用来执行线程束的归约（或子线程束），每一个线程使用不同顺序执行具有结合律的操作计算归约值。

8.6.2 线程块级原语

函数__syncthreads()起着栅栏的作用。它会导致所有的线程等待，直到线程块中的所有线程执行到函数__syncthreads()。费米指令集（SM 2.x）增加了几个新的线程块级的栅栏指令，这些指令均起着聚合线程块中线程信息的作用。

__syncthreads_count(): 评估一个断定并返回断定为真的线程总数。
__syncthreads_or(): 返回线程块中所有线程输入的数值的或运算。
__syncthreads_and(): 返回线程块中所有线程输入的数值的与运算。

8.6.3 性能计数器

开发人员可以定义自己的一套性能计数器，并通过内置函数__prof_trigger()在代码中增加计数值。

void __prof_trigger(int counter);

调用这个函数会使每个线程束中相关的计数器增1。counter必须在范围0~~7之内；计数器8~~15是被保留的。计数器的值可以在配置文件中的prof_trigger_00~prof_trigger_07找到。

8.6.4 视频指令

本节所提到的视频指令只能通过内联PTX汇编代码访问。我们在这里描述它们的基本功能，来帮助开发人员决定视频指令是否对他们的应用有用。任何使用这一指令的人，应该遵循PTX ISA标准。

1. 标量视频指令

标量视频指令，在SM 2.0硬件中添加，实现了在视频处理中在短（8位和16位）整数类型需要的高效操作。按照PTX3.1 ISA规范，这些函数的格式如下。

vop.dtype.atype.btype{.sat}d,a{.asel},b{.bsel}; vop.dtype.atype.btype{.sat}.secop d,a{.asel},b{.bsel},c;

源和目的操作数都是32位的寄存器。dtype、atype和btype可能是.u32或.s32类型，这两种类型分别代表无符号和有符号32位整型数。选择符asel/bsel选择在源操作数中选取哪一个8或16位数：b0、b1、b2和b3选择字节（从低有效位开始计数），h0/h1分别选择最低有效和最高有效的16位。

一旦输入值被提取，通过符号扩展或零扩展为有符号33位整型数，并且基本操作被执行，产生一个34位的中间结果，其符号取决于dtype。最后，其结果夹取到输出范围，并执行以下操作之一：

1）在中间结果和和第三个操作数上应用第二个操作符（add、min或max）。
2）截断中间结果为8或16位值，并合并到第三个操作数指定的位置，以产生最终的结果。

然后写入低32位到目标操作数。

指令vset执行一个在8、16或32位输入操作数之间的比较，并生成相应的断定（1或0）作为输出。PTX标量视频指令和对应操作在表8-14中给出。

表8-14 标量视频指令

2. 矢量视频指令（仅适用于SM 3.0）

在SM3.0中加入的矢量视频指令，类似于标量的视频指令，将输入增强为标准整型格式，执行核心操作，并对输出进行夹取和根据需要进行合并。但通过在成对的16位值或4个一组的8位值上操作，矢量视频指令有着更高的执行性能。

表8-15总结了由这些指令实现的PTX指令和对应操作。它们在视频处理和特定的图像处理中（例如中值过滤）最有用处。

表8-15 矢量视频指令

8.6.5 特殊寄存器

许多特殊寄存器的访问是通过引用内置变量threadIdx、blockIdx、blockDim和gridDim实现的。在7.3节我们详细的描述了这

个三维结构，这些变量分别指定了线程ID、线程块ID、线程数和线程块数。

除了这些伪变量，另一个特殊的寄存器是流处理器簇的时钟寄存器（clock register），它在每个时钟周期增1。这个计数器可以通过函数__clock()或__clock64()读取。时钟计数器独立地跟踪每个流处理器簇，类似于在CPU上的时间戳计数器，在衡量不同代码片段的性能表现时十分有用，但在试图计算挂钟时间时最好避免使用。

8.6_其他指令