目录

背景

介绍

块内组

线程组和线程块

分片（Tiled Partitions）

线程块分片（Thread Block Tiles）

伪线程混洗函数

伪线程投票函数

伪线程匹配函数

合并组

块内合作组的使用

发现模板

伪线程同步代码模板

组合

网格同步

多设备同步

结语

背景

今天我们翻译一下CUDA10.0官方文档中最后一个值得我们注意的部分——合作组

介绍

合作组(Cooperative Groups)是在CUDA 9中为了组织通信线程组而引进的对CUDA编程模型的一个扩展。合作组允许开发者表达线程通信的粒度，以表示更丰富更有效的并行解构。

在这之前(参见文章CUDA编程模型)，CUDA编程模型已经提供了同步合作线程的单一简单结构：__syncthreads()函数里实现的跨线程块内所有线程的栅栏。然而，程序员想要定义并同步其他粒度的线程组的同步，以支持更好的表现、设计灵活性和组级别相关函数接口的软件复用。为了表达更宽泛的并行交互模板，很多面向性能的程序员已经为伪线程内的线程或者单一GPU上的不同线程块的同步实现了自己的、 单独定制但是不安全的函数。尽管达到的性能改善可圈可点，但这导致了零碎代码的泛滥，这种零碎代码随着时间和GPU的更新换代变得难以实现、调整和保持。而合作组通过为支持性能优越的代码提供一个安全且适配未来的机制来解决这一问题。

合作组编程模型扩展在CUDA线程块内和CUDA线程块间都描述了同步模板，这为应用提供了定义他们自己的线程组和同步线程组的方法，也提供了强制使用某些限制的新的启动API，从而保证了同步可以工作。这些函数自吃了CUDA内合作并行的新模板，包括生产者消费者并行、机会主义并行和整个网格内的全局同步。

把组表达为一阶程序对象改善了软件的组成，因为相关函数可以收到表示参与线程组的明确对象，这个对象也让程序员意图明确，那就是消除导致零碎代码、不合理的编译器优化限制的不健全架构假设，并且更好地适配新的GPU版本。

合作组编程模型由下面的元素组成：

• 表示合作线程的数据类型；
• 获取CUDA启动API定义的指令级组的操作；
• 把现存的组分成新组的操作；
• 同步一个给定组的栅栏操作；
• 查看组属性和特定组集合的操作

块内组

本节中我们描述在块内创建能够同步和合作的线程组的可用功能，注意跨线程块或设备的合作组同步需要一些额外的考虑，这在后文会有所描述。

合作组需要CUDA版本>=9.0，为了使用这一功能，需要添加头文件：#include <cooperative_groups.h>，并且使用合作组命名空间：using namespace cooperative_groups;，而后包含任何块内合作组功能的代码可以用nvcc以常规的方式编译。

线程组和线程块

任何CUDA程序员想必都已经对一组线程——线程块(参见文章CUDA编程模型)很熟悉了。合作组拓展引入了一个新的数据类型——thread_block，以在核函数中明确表示这一概念，组可以这样初始化：thread_block g = this_thread_block();，thread_block数据类型来自于更通用的thread_group数据类型。thread_group可以被用来表示范围更宽的组，并提供了以下函数：

void sync(); // 同步组内线程
unsigned size(); // 组内线程数
unsigned thread_rank(); // 调用线程的组内序号，值域为[0, size]
bool is_valid(); // 组是否违反了任何API约束

而thread_block又提供了下面的基于块的额外功能：

dim3 group_index(); // 网格内的块索引，三维
dim3 thread_index(); // 块内的线程索引，三维

比如，如果组g已经按照上面的方式初始化了，那么g.sync();将会同步块内的所有线程，等同于__syncthreads();。注意，组内所有线程都必须执行统一的操作，否则就会产生未定义的行为。

分片（Tiled Partitions）

tile_partition()函数可以用来把线程块解构成多个更小的合作线程组。例如，如果我们先创建了一个包含块内所有线程的组：

thread_block wholeBlock = this_thread_block();

然后，我们可以把它分成更小的组，比如每组32个线程：

thread_group tile32 = tiled_partition(wholeBlock, 32);

更进一步，我们可以把每组32个线程划分成更小的组，比如每组4个线程：

thread_group tile4 = tiled_partition(tile32, 4);

而后，如果我们再加上下面的代码：

if (tile4.thread_rank() == 0) printf(“Hello from tile4 rank 0\n”);

那么，就会每四个线程打印一段话：每个tile4组的第0个线程、同时也是wholeBlock组的第0、4、8、12个线程会进行输出。注意，目前片大小只能是2的整次幂，并且不能超过32

线程块分片（Thread Block Tiles）

tiled_partition函数的模板化版本也可以使用，其中模板参数用来指定分片的大小——其在编译期就确定下来，从而有着更多的执行优化空间。和前一小节的类似，下面的代码将会创建两组分片集合，大小分别是32和4：

thread_block_tile<32> tile32 = tiled_partition<32>(this_thread_block());
thread_block_tile<4> tile4 = tiled_partition<4>(this_thread_block());

注意这里使用了thread_block_tile模板化数据结构，而且组大小是以模板参数而非函数参数传给tiled_partition()函数的。

线程块分片也提供了以下的额外功能：

.shfl()
.shfl_down()
.shfl_up()
.shfl_xor()
.any()
.all()
.ballot()
.match_any()
.match_all()

这些合作同步操作和伪线程混洗函数、伪线程投票函数和伪线程匹配函数类似，这里简单介绍一下。

伪线程混洗函数

伪线程混洗函数用来在不使用共享内存的情况下实现在伪线程内的线程之间广播数据，函数原型如下：

T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width=warpSize);
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int
width=warpSize);
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize);

其中T为要广播的数据类型，可以是int、unsigned int、long、unsigned long、long long、unsigned long long、float或double，如果包含了头文件cuda_fp16.h，T也可以是__half或__half2；mask用来标记执行交换的目标线程；srcLane表示发送广播的源线程，如果源线程id大于width，那么实际的源线程id就等于srcLane % width；width表示执行广播的分组大小，必须是2的整次幂，而且不超过32，在指定大小的组内会进行值的广播；函数返回的是源线程中value指定的四字节的字；

__shfl_sync()函数源线程id就是srcLane；__shfl_up_sync()函数的源线程id为srcLane - delta；__shfl_down_sync()函数的源线程id为srcLane + delta；__shfl_xor_sync()函数的源线程id为srcLane xor laneMask

伪线程投票函数

伪线程投票函数允许伪线程内的线程执行归约广播操作，这些函数的原型如下：

int __all_sync(unsigned mask, int predicate);
int __any_sync(unsigned mask, int predicate);
unsigned __ballot_sync(unsigned mask, int predicate);
unsigned __activemask();

predicate表示判断谓语，mask表示参与投票的线程。函数从伪线程中的每个线程里读取整型的谓语，并这些谓语值是否为0，并把返回值广播给每个参与的线程。函数的执行逻辑如下表所示

函数	执行逻辑
__all_sync()	给mask指定的所有未退出线程评估谓语值，只有所有线程的谓语值均非零，才返回非零值
__any_sync()	给mask指定的所有未退出线程评估谓语值，任一线程的谓语值非零，就返回非零值
__ballot_sync()	给mask指定的所有未退出线程评估谓语值，返回一个整数。当且仅当伪线程的第N个线程是活动的、且谓语值非0，整数的第N位才为1
__activemask()	返回伪线程中目前所有的活动线程的4字节掩码。如果调用此函数时伪线程中的第N个线程是活动的，掩码的第N位就为1，已退出或非活动线程对应的码位为0。注意，此函数调用时收敛的线程不能保证在下游指令处依旧是收敛的，除非这些指令是伪线程内置的同步函数

伪线程匹配函数

伪线程匹配函数会在伪线程中线程间执行同步后的广播比较操作，支持计算能力>=7.X的设备，函数原型如下：

unsigned int __match_any_sync(unsigned mask, T value);
unsigned int __match_all_sync(unsigned mask, T value, int *pred);

T可以是int、unsigned int、long、unsigned long、long long、unsigned long long、float或double，value表示要广播比较的值，mask指定要参与的线程。这两个函数的返回逻辑不同，如下表所示

函数	返回逻辑
__match_any_sync()	返回mask指定的线程中拥有和value相等的值的线程掩码
__match_all_sync()	如果mask指定的线程都拥有和value相等的值，才返回mask，同时pred为真；否则返回0，pred为假

回到合作组线程的线程块分片小节中，其中的这些函数是在用户定义线程组的上下文中使用的，并且提供了更好的灵活性和生产效率。

合并组

在CUDA的SIMT架构(参见文章CUDA硬件实现)中，在硬件层面，多处理器以32个线程为一组(伪线程)来执行线程。如果在应用代码中存在数据依赖的条件分支而导致的伪线程内的线程分散，那么伪线程就会穿行执行每个分支，同时阻塞不在那条路径上的线程，而在当前执行路径上的活跃线程的执行称之为合并执行。合作组有发现或创建包含所有合并线程组的功能：coalesced_group active = coalesced_threads();。例如，考虑一个场景：代码中存在着只有每个伪线程中的第2、4、8个线程保持活跃的分支，在此分支中执行方才的那条语句就会为每个伪线程创建名为active的组，包含那三个活跃的线程(组内id分别为0、1、2)

块内合作组的使用

在本小节中，合作组的功能通过一些例子来阐述

发现模板

一般情况下，开发者需要和活跃线程集工作，我们不能假设或指定当前有哪些线程，而只能和碰巧处于活动的线程(threads that happen to be there)工作，这一点可以从下面的“伪线程内跨线程的聚集原子加”中看到（使用正确的CUDA 9.0函数写的）：

{unsigned int writemask = __activemask();unsigned int total = __popc(writemask); // 活跃的线程数unsigned int prefix = __popc(writemask & __lanemask_lt()); // 当前活跃线程前缀，比如活跃线程掩码为01010，那么对于第2个线程，__lanemask_lt()为00001，那么prefix就是0(第4个活跃线程对应的就是1)。因此前缀为0就表示当前为第一个活跃的线程int elected_lane = __ffs(writemask) - 1; // id最小的活跃线程int base_offset = 0;if (prefix == 0) {base_offset = atomicAdd(p, total);}base_offset = __shfl_sync(writemask, base_offset, elected_lane); // 把elected_lane中原子加前的值广播到所有的活跃线程中int thread_offset = prefix + base_offset;return thread_offset;
}

如果用合作组API重写，就会得到以下代码：

{cg::coalesced_group g = cg::coalesced_threads(); // 活跃线程组int prev;if (g.thread_rank() == 0) { // 第一个活跃线程prev = atomicAdd(p, g.size()); // 原子加}prev = g.thread_rank() + g.shfl(prev, 0); // 最小的活跃线程id + 老的值return prev;
}

伪线程同步代码模板

开发者可能有过伪线程同步的代码，并且做出了伪线程大小的隐式假设，并且根据这个大小进行编码。现在，伪线程大小需要被明确指定：

auto g = tiled_partition<16>(this_thread_block());

然而，用户可能想更好地对算法进行分区，而且不使用伪线程同步内置模板参数：

auto g = tiled_partition(this_thread_block(), 8);

在这种情况下，组g依旧可以同步，而且我们依旧可以基于其来构造多种并行算法，但是shfl()等函数就不能使用了：

__global__ void cooperative_kernel(...) {// 获取默认的块线程组thread_group my_block = this_thread_block();// 分组成32个线程一组的线程组(片)，线程片将线程组平均瓜分，每个片内的线程都是连续的thread_group my_tile = tiled_partition(my_block, 32);// 只在块内前32个线程中执行操作if (my_block.thread_rank() < 32) {// ...my_tile.sync();}
}

组合

以前，写代码时，对实现有着一些隐式限制，比如以下代码：

__device__ int sum(int *x, int n) {// ...__syncthreads();return total;
}__global__ void parallel_kernel(float *x){// ...// 所有的线程块都要调用sum()sum(x, n);
}

线程块内的线程必须到达__syncthreads()栅栏，但是这个限制对调用sum()的开发者是不可见的。那么，使用合作组，更好的实现方式可以是：

__device__ int sum(const thread_group& g, int *x, int n)
{// ...g.sync()return total;
}__global__ void parallel_kernel(...)
{// ...sum(this_thread_block(), x, n);// ...
}

网格同步

在引入合作组同步之前，CUDA编程模型只允许核函数完成时的线程块之间的同步，核函数界限有着隐式地无效状态和潜在的性能影响。例如，在特定的用例下，应用有大量的小核函数，每个核函数表示流水线的一个阶段。目前的CUDA编程模型要求这些核函数在操作下一流水线阶段的线程块准备好消费数据前，操作当前流水线阶段的线程块就要生产出数据。在这种情况下，提供全局线程块间同步的能力将会允许应用重构这些线程块，以在一个给定阶段完成时同步设备。

为了在一个核函数内部同步网格，可以使用组：grid_group grid = this_grid();，然后调用grid.sync();。为了支持单元格同步，当启动核函数时我们有必要使用cudaLaunchCooperativeKernel()这一CUDA运行时启动API，而不是<<<>>>执行配置语法：

cudaLaunchCooperativeKernel(const T *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem = 0, cudaStream_t stream = 0)
// 或者CUDA驱动API的对应函数，这种核函数不能使用附录A中的动态并行功能

为了确保线程块在GPU上的共存性，启动的块数量需要被小心考量。比如，我们可以按照下面的方式启动：

cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
// 初始化，而后启动
cudaLaunchCooperativeKernel((void*)my_kernel, deviceProp.multiProcessorCount, numThreads, args);

或者，我们可以按照下面的方式使用占有率计算器来计算多少线程块可以同时存在于一个多处理器上：

cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, my_kernel, numThreads, 0));
// 初始化，而后启动
cudaLaunchCooperativeKernel((void*)my_kernel, numBlocksPerSm, numThreads, args);

也要注意，为了使用网格同步，设备代码必须被单独编译，然后设备运行时再链接进入(详情请参见CUDA Compiler Driver NVCC文档的在CUDA中使用独立编译章节)，最简单的例子如下所示：

nvcc -arch=sm_61 -rdc=true mytestfile.cu -o mytest

我们也要确保设备支持合作启动属性，可以使用cuDeviceAttribute()这一CUDA驱动API来查看：

int pi=0;
cuDevice dev;
cuDeviceGet(&dev,0) // 查询设备0
cuDeviceGetAttribute(&pi, CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH, dev);

如果pi为1，就说明这一属性在设备0上支持，只有计算能力≥6.0的设备才能支持合作启动属性。另外，我们应该在没有MPS的Linux平台或者使用TCC模式设备的Winwos平台运行带有合作启动功能的程序

多设备同步

为了支持多个使用合作组设备之间的同步，我们需要使用cuLaunchCooperativeKernelMultiDevice()这一CUDA api，这是对现有CUDA API的一个重要扩展，将会支持一个主机线程在多个设备上启动核函数。除了cuLaunchCooperativeKernel()函数之外的限制和保证，cuLaunchCooperativeKernelMultiDevice()函数还有着以下的语义：

• 这个API保证启动是原子的，也就是说如果API成功调用，那么所有指定的设备上将启动指定数量的线程块；
• 通过此API启动的核函数必须相同，这部分驱动不会做显式的检查，因为这种检查在驱动中基本不可行，所以应该让应用保证这一点；
• launchParamsList参数中的两个元素不能映射到一台设备上；
• 这种启动的目标设备必须有相同的计算能力——主版本或副版本都得相等；
• 线程块大小、网格大小和每个单元格使用的共享内存数必须对于所有设备都相等。注意，这意味着每台设备启动的最大线程块数取决于拥有最少多处理器数量的设备；
• 任何存在于调用待启动cu函数的模块里的自定义__device__、__constant__或__managed__设备全局变量会在每台设备上独立地初始化，用户应该保证这种设备全局变量初始化的正确性。

启动参数应该用下面的结构体定义：

typedef struct CUDA_LAUNCH_PARAMS_st {CUfunction function;unsigned int gridDimX;unsigned int gridDimY;unsigned int gridDimZ;unsigned int blockDimX;unsigned int blockDimY;unsigned int blockDimZ;unsigned int sharedMemBytes;CUstream hStream;void **kernelParams;
} CUDA_LAUNCH_PARAMS;

然后传给启动API：

cudaLaunchCooperativeKernelMultiDevice(CUDA_LAUNCH_PARAMS *launchParamsList, unsigned int numDevices, unsigned int flags = 0);

这种启动方式和上面网格同步的启动类似，同样类似的还有同步方式：

multi_grid_group multi_grid = this_multi_grid();
multi_grid.sync();

也需要使用独立编译。

我们应该确保设备支持多设备启动属性，方式和前一节描述的类似，只需要把参数换成CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH，只有计算能力≥6.0的设备才能支持合作启动属性。另外，我们应该在没有MPS的Linux平台或者使用TCC模式设备的Winwos平台运行带有合作启动功能的程序

结语

到此，我对CUDA10.0官方文档的翻译就分享完了，全程亲自翻译，但英语水平有限，如有不当之处，还请在评论区提出建议，劳驾。