1. 单指令多线程模式

从硬件上看，一个GPU被分为若干个SM。线程块在执行时将被分配到还没完全占满的SM中，一个线程块不会被分配到不同的SM中，一个SM可以有一个或多个线程块。不同线程块之间可以并发或顺序地执行。当某些线程块完成计算任务后，对应的SM会部分或完全地空闲，然后会有新的线程块被分配到空闲的SM。从更细的粒度看，一个SM以32个线程为单位产生、管理、调度、执行线程，这样的32个线程称为一个线程束，每个线程束包含32个具有连续线程号的线程。

在Volta架构之前，一个线程束中的线程拥有同一个程序计数器(program counter)，但有各自不同的寄存器状态。在同一时刻，一个线程束中的线程只能执行一个共同的指令或者闲置，这称为单指令多线程(single instruction multiple thread, SIMT)模式。当一个线程束中的线程顺序地执行判断语句中的不同分支时，即发生了分支发散(branch divergence)。例如有如下语句：

if (condition)
{
    A;
}
else
{
    B;
}

首先，满足condition条件的线程会执行语句A，其它线程闲置；然后，不满足condition条件的线程执行语句B，其它线程闲置。如果语句A和语句B的指令数差不多，则整个线程束的执行效率就会降低一半，所以在编写代码时，应该尽量避免分支发散。需要注意的是，分支发散是针对同一个线程束内部线程的，不同线程束执行条件语句的不同分支则不属于分支发散。

从Volta架构开始，引入了独立线程调度机制，每个线程拥有自己的程序计数器。同时，这也使得假设了线程束同步的代码变得不再安全。如果要在Volta或者更高架构的GPU中运行一个使用了线程束同步假设的程序，可以在编译时将虚拟架构指定为低于Volta架构的计算能力，例如-arch=compute_60 -code=sm_70，这样在生成PTX代码时就使用了Pascal架构的线程调度机制，而忽略了Volta架构的独立线程调度机制。

2. 线程同步

线程块同步函数：

//保证一个线程块中的所有线程（或者说所有线程束）在执行该语句后面的语句之前都执行完了该语句前面的语句
void __syncthreads();

线程束同步函数：

//参数mask是一个代表掩码的无符号整数，默认32个比特位都为1，表示线程束中的所有线程都参与同步，如果要排除一些线程，可以把对应比特位置0，例如0xfffffffe表示排除第0号线程
void __syncwarp(unsigned mask=0xffffffff);

此外，还有一些线程束内的基本函数，它们都具有隐式的同步功能。其中线程束表决函数(warp vote functions)和线程束洗牌函数(warp shuffle functions)自Kepler架构开始就可以使用，但在CUDA 9版本中进行了更新，线程束匹配函数(warp match functions)和线程束矩阵函数(warp matrix functions)只能在Volta或更高架构的GPU中使用。

线程束内基本函数中的参数mask称为掩码，是一个32位的无符号整数，其二进制从右边数起刚好对应线程束内的32个线程。掩码用于指定要参与计算的线程，比特位等于1表示参与计算，比特位等于0表示忽略。各种函数返回的结果对于被掩码排除的线程来说没有定义，所以不要在被排除的线程中使用函数的返回值。

// ================ 线程束表决函数 ================

//如果线程束内第n个线程参与计算且pred值非0，则返回无符号整数的第n个比特位取1，否则取0。该函数相当于从一个旧的掩码产生一个新的掩码。
unsigned __ballot_sync(unsigned mask, int pred);

//线程束内所有参与线程的pred值都不为0时才返回1，否则返回0。该函数类似于这样一种选举操作，当所有参选人都同意时才通过。
int __all_sync(unsigned mask, int pred);

//线程束内所有参与线程的pred值至少有一个不为0时就返回1，否则返回0。该函数类似于这样一种选举操作，只要有一个参选人同意就通过。
int __any_sync(unsigned mask, int pred);

// ================ 线程束洗牌函数 ================
//对于所有洗牌函数，类型T可以是int、long、long long、unsigned、unsigned long、unsigned long long、float、double。
//最后一个参数width默认值为warpSize（即32），且只能取2、4、8、16、32其中的一个，它表示逻辑上线程束的大小。
//标号srcLane指的是当前线程在width范围内的位置，例如当width等于8时，srcLane的范围就是0~7。

//参与线程返回标号为srcLane的线程中变量var的值，即将一个线程的数据广播到所有（包括自己）线程。
T __shfl_sync(unsigned mask, T var, int srcLane, int width);

//标号为srcLane的参与线程返回标号为srcLane - delta的线程中变量var的值，标号srcLane < delta的线程返回自己的var值。形象地说，这是一种将数据向上平移的操作。
T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width);

//标号为srcLane的参与线程返回标号为srcLane + delta的线程中变量var的值，标号srcLane >= width - delta的线程返回自己的var值。形象地说，这是一种将数据向下平移的操作。
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width);

//标号为srcLane的参与线程返回标号为srcLane ^ laneMask的线程中变量var的值，这里的^符号表示整数按位异或的操作。例如width等于8，laneMask等于2时，第0~7号线程分别返回标号为2、3、0、1、6、7、4、5的线程中变量var的值。
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width)

为了更好地理解上述函数，可以参考下面的测试程序，输出结果在注释中：

#include <cstdio>

const unsigned WIDTH = 8;
const unsigned BLOCK_SIZE = 16;
const unsigned FULL_MASK = 0xffffffff;

void __global__ test_warp_primitives(void);

int main(int argc, char** argv)
{
    test_warp_primitives<<<1, BLOCK_SIZE>>>();
    cudaDeviceSynchronize();
    return 0