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