并行编程实战——CUDA编程的协作组

一、协作组（Cooperative Groups ）

Cooperative Groups,协作组。大家可能对不断出现各种新的名词感到很困惑，其实原因非常简单，在实际的应用中会出现各种各样的需求，而为了达到解决某种需求的目的，实现的一种技术就会产生一种新名词。就像大家非常熟悉汽车，但对汽车行业提出的各种技术名词仍然不了解一样。
从CUDA9开始，允许开发者通过协作组来实现一种新规格的线程通信粒度的表现形式，可以更加有效的对并行进行分解。说得更加直白一些，就是更加容易对并行任务的细节进行了控制。
其实对于一些特定开发者来说，他们早就已经不满足于CUDA提供的线程控制粒度。这就和下馆子一样，同一个人，同样的饭食，可能今天满意，明天就不满意了。不断的变化的需求，才是进步的根本。
大家可以把协作组当成一种线程管理的抽象，至于能够抽象到何种地步，这和CUDA官方提供的技术支撑为基础。比如可以跨Warp和Block，甚至是Grid。从CUDA9开始，到最新的版本中，协作组在不断的演进，比如在CUDA13中移除了multi_grid_group。所以大家在使用一些具体的接口时，要紧密的和硬件及CUDA版本匹配，否则就可能会出现一些低级的错误。
CUDA中的协作组可以分为两类：

1. 隐式
   它包括块协作组（thread_block）;簇协作组（cluster_group）以及网格协作组（grid_group）
2. 显式
   它包括线程块切片（thread_block_tile）和合并协作组（coalesced_group）

二、协作组的分析说明

通过上面的介绍，可以明白了协作组是一种线程编程模型中的同步模式，它从底层提供了自定义线程组的方法和同步的接口。通过一系列的相关接口，可以利用原语保证CUDA线程内部协作并行的相关机制。从而达到在Warp、Block及Grid间的全局同步。它主要包括以下几部分：

1. 协作组的数据类型
2. 获得由CUDA API启动后定义的隐式组的操作，如线程块
3. 将现有组划分为新组的集合
4. 数据移动和操作的集合算法如memcpy_async， reduce, scan
5. 同步组内所有线程的操作
6. 检查线程组属性的操作
7. 暴露低级、特定组的、通常是硬件加速的操作的集合

到目前为止，大家可以明白了线程组是什么了。它就是抽象的线程集，它更利用从编译器到GPU运行时的整体的管理和操作。用面向对象编程的方式来看，它就是一个线程组的对象。在CUDA的文档说明中，提到了为什么不使用泛型而使用线程组，因为前者为降低编译器的优化程度。

三、协作组划分和组操作

协作组可以有三种形式进行分区操作：

1. tiled_partition
   这个有点类似于一维的数组，然后被指定大小进行平均分配划分成多个子组。
2. labeled_partition
   就是用一种标记划分各个子组，通过标签来协调线程的同步
3. binary_partition
   它是labeled_partition划分的一种特殊形式，通过谓词计算得到的结果来划分组，但其标签只有0和1两种情况。

划分组的目的是什么？当然是进行各种操作处理。它主要有：

1. 同步
   通过barrier_arrive 、barrier_wait和sync接口实现了线程间的同步
2. 数据交换
   通过memcpy_async和wait、wait_prior实现数据的操作
3. 数据控制
   利用reduce等接口实现数据的扫描和减少操作
4. 运行控制
   就是任务的调用和分发机制

四、协作组的应用

下面看一个简单例程：

#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <cooperative_groups.h>
#include <stdio.h>
#include <iostream>

using namespace cooperative_groups;

__global__ void blockGroup(int* data) {
	// 获取当前线程块组
	thread_block block = this_thread_block();

	int tid = block.thread_rank();
	int bid = block.group_index().x;

	// 块内同步
	data[bid * blockDim.x + tid] = tid * 2;

	block.sync();  // 线程块内所有线程同步

	//同步后继续工作
	if (tid == 0) {
		data[bid * blockDim.x] *= 100;
	}
}


void printData(int* arr, int size, const char* name) {
	printf("%s: ", name);
	for (int i = 0; i < size; i++) {
		printf("%d ", arr[i]);
	}
	printf("\n");
}

int main() {
	const int N = 256;
	const int BLOCK_SIZE = 128;
	const int GRID_SIZE = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;

	int * dData;
	int * hData;

	// 分配主机内存
	hData = new int[N];

	cudaMalloc(&dData, N * sizeof(int));

	// 初始化
	int* hiData = new int[N];
	for (int i = 0; i < N; i++) {
		hiData[i] = ++i;
	}

	cudaMemcpy(dData, hiData, N * sizeof(int), cudaMemcpyHostToDevice);

	printf("\n----------start test--------------:\n");
	blockGroup << <GRID_SIZE, BLOCK_SIZE >> >(dData);
	cudaGetLastError();
	cudaMemcpy(hData, dData, N * sizeof(int), cudaMemcpyDeviceToHost);
	printData(hData, 20, "只显示前20个结果");

	delete[] hData;
	delete[] hiData;

	cudaFree(dData);

	printf("\n  ----------END--------------------\n");

	cudaDeviceReset();
	return 0;
}

运行结果是：

----------start test--------------:
只显示前20个结果: 0 2 4 6 8 10 12 14 16 18 20 22 24 26 28 30 32 34 36 38
 ----------END--------------------

五、总结

协作组的出现是一种必然的结果，毕竟实际的需求中，对线程管理的粒度是分层且丰富的。粗放的管理只适合于技术的初始期，一旦被应用开来，各种实地的场景的要求会纷至沓来。这也是为什么一种技术，一个语言不断的膨胀的原因。
其实回到那句老话“创业难，守业更难”。道理都是相通的啊。