《CUDA编程：基础与实践》读书笔记(3)：同步、协作组、原子函数

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;
}

void __global__ test_warp_primitives(void)
{
    int tid = threadIdx.x;
    int lane_id = tid % WIDTH;

    if (tid == 0) printf("threadIdx.x: ");
    printf("%2d ", tid);
    if (tid == 0) printf("\n");
    // threadIdx.x:  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15

    if (tid == 0) printf("lane_id:     ");
    printf("%2d ", lane_id);
    if (tid == 0) printf("\n");
    // lane_id:      0  1  2  3  4  5  6  7  0  1  2  3  4  5  6  7

    unsigned mask1 = __ballot_sync(FULL_MASK, tid > 0);
    unsigned mask2 = __ballot_sync(FULL_MASK, tid == 0);
    if (tid == 0) printf("FULL_MASK = %x\n", FULL_MASK);
    if (tid == 1) printf("mask1     = %x\n", mask1);
    if (tid == 0) printf("mask2     = %x\n", mask2);
    // FULL_MASK = ffffffff
    // mask1     = fffe
    // mask2     = 1

    int result = __all_sync(FULL_MASK, tid);
    if (tid == 0) printf("all_sync (FULL_MASK): %d\n", result);
    // all_sync (FULL_MASK): 0

    result = __all_sync(mask1, tid);
    if (tid == 1) printf("all_sync     (mask1): %d\n", result);
    // all_sync     (mask1): 1

    result = __any_sync(FULL_MASK, tid);
    if (tid == 0) printf("any_sync (FULL_MASK): %d\n", result);
    // any_sync (FULL_MASK): 1

    result = __any_sync(mask2, tid);
    if (tid == 0) printf("any_sync     (mask2): %d\n", result);
    // any_sync     (mask2): 0

    int value = __shfl_sync(FULL_MASK, tid, 2, WIDTH);
    if (tid == 0) printf("shfl:      ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");
    // shfl:       2  2  2  2  2  2  2  2 10 10 10 10 10 10 10 10

    value = __shfl_up_sync(FULL_MASK, tid, 1, WIDTH);
    if (tid == 0) printf("shfl_up:   ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");
    // shfl_up:    0  0  1  2  3  4  5  6  8  8  9 10 11 12 13 14

    value = __shfl_down_sync(FULL_MASK, tid, 1, WIDTH);
    if (tid == 0) printf("shfl_down: ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");
    // shfl_down:  1  2  3  4  5  6  7  7  9 10 11 12 13 14 15 15

    value = __shfl_xor_sync(FULL_MASK, tid, 1, WIDTH);
    if (tid == 0) printf("shfl_xor:  ");
    printf("%2d ", value);
    if (tid == 0) printf("\n");
    // shfl_xor:   1  0  3  2  5  4  7  6  9  8 11 10 13 12 15 14
}

3. 协作组

协作组(cooperative groups)可以看作线程块和线程束同步机制的推广，它提供了更为灵活的线程协作方式，包括线程块内部的同步协作、线程块之间（网格级）的同步协作以及设备之间的同步协作，本文只介绍线程块内部的协作组。协作组由CUDA 9引入，使用协作组需要包含对应头文件并使用命名空间：

#include "cooperative_groups.h"

using namespace cooperative_groups;

协作组编程模型中最基本的类型是线程组thread_group，它的成员如下：

//同步组内所有线程
void sync();

//返回组内总的线程数目
unsigned int size();

//返回当前调用该函数的线程在组内的标号（从0开始计数）
unsigned int thread_rank();

//如果定义的组违反了任何CUDA限制则返回false，否则返回true
bool is_valid();

thread_block派生自thread_group，并提供了额外的函数：

//返回当前调用该函数的线程的线程块标号，等价于blockIdx
dim3 group_index();

//返回当前调用该函数的线程的线程标号，等价于threadIdx
dim3 thread_index();

可以用如下方式定义并初始化一个thread_block对象：

thread_block tb = this_thread_block();

其中this_thread_block()相当于一个线程块类型的常量，这样定义的tb就代表当前线程块，只不过这里把它包装成了一个类型。例如，tb.sync()完全等价于__syncthreads()，tb.group_index()完全等价于blockIdx，tb.thread_index()完全等价于threadIdx。

可以用函数tiled_partition将一个线程块划分成若干片(tile)，每片构成一个新的线程组，目前仅可将片的大小设置为2、4、8、16、32中的一个。线程组也可以被分割为更细的线程组。

thread_group g32 = tiled_partition(this_thread_block(), 32);
thread_group g4 = tiled_partition(g32, 4);

如果线程组的大小在编译期就已知，那么就可以使用模板化的版本进行定义，这样可能会更高效。

thread_block_tile<32> g32 = tiled_partition<32>(this_thread_block());

对于用模板定义的线程块片(thread block tile)，还可以使用如下函数（类似线程束内的基本函数）：

int any(int predicate);
int all(int predicate);
unsigned int ballot(int predicate);
T shfl(T var, int srcLane);
T shfl_down(T var, unsigned int delta);
T shfl_up(T var, unsigned int delta);
T shfl_xor(T var, unsigned int laneMask);

与对应的线程束内基本函数相比，上述线程块片的函数主要有两点不同：①、少了代表掩码的参数，因为线程组内所有线程都必须参与函数的计算；②、洗牌函数少了最后一个代表宽度的参数，因为宽度就等于线程块片的大小，即模板参数。

4. 原子函数

原子函数对它第一个参数指向的数据进行一次“读-改-写”的原子操作，第一个参数可以指向全局内存，也可以指向共享内存。原子操作是一个线程一个线程轮流做的，但没有明确的次序，另外，原子函数没有同步功能。所有原子函数都是__device__函数，只能在核函数中使用。

下表列出了所有原子函数的原型，address所指变量的值在执行原子函数前为old，执行原子函数后为new。对于每个原子函数，返回值都是old。

运算	功能	函数原型
加法	new = old + val	T atomicAdd(T* address, T val);
减法	new = old - val	T atomicSub(T* address, T val);
自增	new = (old >= val) ? 0 : (old + 1)	T atomicInc(T* address, T val);
自减	new = ((old == 0) || (old > val)) ? val : (old - 1)	T atomicDec(T* address, T val);
最小值	new = (old < val) : old : val	T atomicMin(T* address, T val);
最大值	new = (old > val) : old : val	T atomicMax(T* address, T val);
交换	new = val	T atomicExch(T* address, T val);
比较-交换(compare and swap)	new = (old == compare) ? val : old	T atomicCAS(T* address, T compare, T val);
按位与	new = old & val	T atomicAnd(T* address, T val);
按位或	new = old | val	T atomicOr(T* address, T val);
按位异或	new = old ^ val	T atomicXor(T* address, T val);

上面所列的函数中，我们用T表示相关变量的数据类型。各个原子函数对数据类型的支持情况见下表。注：atomicAdd对double和__half2的支持始于Pascal架构，对__half的支持始于Volta架构。

原子函数	int	unsigned	unsigned long long	float	double	__half2	__half
atomicAdd	yes	yes	yes	yes	yes	yes	yes
atomicSub	yes	yes	no	no	no	no	no
atomicInc	no	yes	no	no	no	no	no
atomicDec	no	yes	no	no	no	no	no
atomicMin	yes	yes	yes	no	no	no	no
atomicMax	yes	yes	yes	no	no	no	no
atomicExch	yes	yes	yes	yes	no	no	no
atomicCAS	yes	yes	yes	no	no	no	no
atomicAnd	yes	yes	yes	no	no	no	no
atomicOr	yes	yes	yes	no	no	no	no
atomicXor	yes	yes	yes	no	no	no	no