线程束（warp），即一个线程块中连续32个线程。
一个GPU被分为若干个流多处理器（SM）。核函数中定义的线程块（block）在执行时
将被分配到还没有完全占满的 SM。
一个block不会被分配到不同的SM，同时一个 SM 中可以有多个 block。不同的block
之间可以并发也可以顺序执行，一般不能同步。当某些block完成计算任务后，对应的
SM 会部分或完全空闲，然后会有新的block被分配到空闲的SM。
一个 SM 以32个线程（warp）为单位产生、管理、调度、执行线程。
一个 SM 可以处理多个block，一个block可以分为若干个warp。
在同一时刻，一个warp中的线程只能执行一个共同的指令或者闲置，即单指令-多线程执行模型，
（single instruction multiple thread, SIMT）。
当一个线程束中线程顺序的执行判断语句中的不同分支时，称为发生了 分支发散（branch divergence）。
if (condition)
{
A;
}
else
{
B;
}
首先，满足 condition 的线程或执行语句A，其他的线程会闲置；
然后，不满足条件的将会执行语句B，其他线程闲置。
当语句A和B的指令数差不多时，整个warp的执行效率就比没有分支的情况 低一半。
一般应当在核函数中尽量避免分支发散，但有时这也是不可避免的。
如数组计算中常用的判断语句：
if(n < N)
{
// do something.
}
该分支判断最多影响最后一个block中的某些warp发生分支发散，
一般不会显著地影响性能。
有时能通过 合并判断语句 的方式减少分支发散；另外，如果两分支中有一个分支
不包含指令，则即使发生分支发散也不会显著影响性能。
注意不同架构中的线程调度机制
__syncwarp(·)：当所涉及的线程都在一个线程束内时，可以将线程块同步函数 __syncthreads()
换成一个更加廉价的线程束同步函数__syncwarp(·)，简称 束内同步函数。
函数参数是一个代表掩码的无符号整型数，默认值是全部32个二进制位都为1，代表
线程束中的所有线程都参与同步。
关于 掩码(mask) 的简介文章：思否小姐姐：“奇怪的知识——位掩码”
线程束表决函数；
-
unsigned __ballot_sync(unsigned mask, int predicate)，
如果线程束内第n个线程参与计算（旧掩码）且predicate值非零，则返回的无符号整型数（新掩码）
的第n个二进制位为1，否则为0。 -
int __all_sync(unsigned mask, int predicate)，
线程束内所有参与线程的predicate值均非零，则返回1，否则返回0. -
int __any_sync(unsigned mask, int predicate)，
线程束内所有参与线程的predicate值存在非零，则返回1， 否则返回0.
线程束洗牌函数：
-
T __shfl_sync(unsigned mask, T v, int srcLane, int w = warpSize)，
参与线程返回标号为 srcLane 的线程中变量 v 的值。
该函数将一个线程中的数据广播到所有线程。 -
T __shfl_up_sync(unsigned mask, T v, unsigned d, int w=warpSize)，
标号为t的参与线程返回标号为 t-d 的线程中变量v的值，t-d<0的线程返回t线程的变量v。
该函数是一种将数据向上平移的操作，即将低线程号的值平移到高线程号。
例如当w=8、d=2时，2-7号线程将返回 0-5号线程中变量v的值；0-1号线程返回自己的 v。 -
T __shfl_down_sync(unsigned mask, T v, unsigned d, int w=warpSize)，
标号为t的参与线程返回标号为 t+d 的线程中变量v的值，t+d>w的线程返回t线程的变量v。
该函数是一种将数据向下平移的操作，即将高线程号的值平移到低线程号。
例如当w=8、d=2时，0-5号线程将返回2-7号线程中变量v的值，6-7号线程将返回自己的 v。 -
T __shfl__xor_sync(unsigned mask, T v, int laneMask, int w=warpSize),
标号为t的参与线程返回标号为 t^laneMask 的线程中变量 v 的值。
该函数让线程束内的线程两两交换数据。
每个线程束洗牌函数都有一个可选参数 w，默认是线程束大小（32），且只能取2、4，8、16、32。
当 w 小于 32 时，相当于逻辑上的线程束大小是 w，其他规则不变。
此时，可以定义一个 束内索引：(假设使用一维线程块)
int laneId = threadIdx.x % w; // 线程索引与束内索引的对应关系。
假设线程块大小为16，w 为 8：
线程索引： 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
束内索引： 0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7
参数中的 mask 称为掩码，是一个无符号整型，具有32位，一般用 十六进制 表示：
const unsigned FULL_MASK = 0xffffffff; // `0x`表示十六进制数；`0b`表示二进制数。
或者
#define FULL_MASK 0xffffffff
以上所有线程束内函数都有 _sync 后缀，表示这些函数都具有 隐式的同步功能。
协作组（cooperative groups），可以看作是线程块和线程束同步机制的推广，
提供包括线程块内部的同步与协作、线程块之间（网格级）的同步与协作、以及
设备与设备之间的同步与协作。
使用协作组需要包含如下头文件：
#include <cooperative_groups.h>
using namespace cooperative_groups;
线程块级别的协作组
协作组编程模型中最基本的类型是线程组 thread_group，其包含如下成员：
void sync()，同步组内所有线程；unsigned size()，返回组内总的线程数目，即组的大小；unsigned thread_rank()，返回当前调用该函数的线程在组内的标号（从0计数）；bool is_valid()，如果定义的组违反了任何cuda限制，返回 false，否则true；
线程组类型有一个导出类型，线程块thread_block，其中定义了额外的函数：
dim3 group_index()，返回当前调用该函数的线程的线程块指标，等价于blockIdx；dim3 thread_index()，返回当前调用该函数的线程的线程指标，等价于threadIdx；
通过 this_thread_block() 初始化一个线程块对象：
thread_block g = this_thread_block(); // 相当于一个线程块类型的常量。
此时，
g.sync() <===> __syncthreads()
g.group_index() <===> blockIdx
g.thread_index() <===> threadIdx
通过 tiled_partition() 可以将一个线程块划分为若干片（tile），每一片构成一个新的线程组。
目前，仅支持将片的大小设置为 2 的整数次方且不大于 32。
thread_group g32 = tiled_partition(this_thread_block(), 32); // 将线程块划分为线程束。
可以继续将线程组划分为更细的线程组：
thread_group g4 = tiled_partition(g32, 4);
采用模板、在编译期划分 线程块片（thread block tile）：
thread_block_tile<32> g32 = tiled_partition<32>(this_thread_block());
thread_block_tile<32> g4 = tiled_partition<4>(this_thread_block());
线程块片具有额外的函数（类似线程束内函数）：
unsigned ballot(int predicate);int all(int predicate);int any(int predicate);T shfl(T v, int srcLane);T shfl_up(T v, unsigned d);T shfl_down(T v, unsigned d);T shfl_xor(T v, unsigned d);
与一般的线程束不同，线程组内的所有线程都要参与代码运行计算；同时，线程组内函数不需要指定宽度，
因为该宽度就是线程块片的大小。
提高线程利用率
在当前的归约程序中，当 offset=64，只用了 1/2 的线程；当 offset=32，只用了 1/4 的线程；...
最终，当 offset=1，只用了 1/128 的线程；
归约过程一共用了 log2(128) = 7 步，平均线程利用率 (1/2 + 1/4 + ... + 1/128)/7 => 1/7。
而在归约前的数据拷贝中线程利用率为 100%，可以尽量把计算放在在归约前：让一个线程处理多个数据。
一个线程处理相邻若干个数据会导致全局内存的非合并访问。要保证全局内存的合并访问，这里需要
保证相邻线程处理相邻数据，一个线程访问的数据需要有某种跨度。
该跨度可以是线程块的大小，也可以是网格的大小；对于一维情况，分别是 blockDim.x 和 blockDim.x * gridDim.x。
避免反复分配与释放设备内存
设备内存的分配与释放是比较耗时的。
通过采用静态全局内存替代动态全局内存，实现编译期的设备内存分配可以更加高效。
此外，应当尽量避免在较内存循环反复的分配和释放设备内存。