4.4. 协作组

4.4.3.1. 尽早创建隐式组句柄

为了获得最佳性能，建议您预先创建隐式组的句柄（尽可能早，在任何分支发生之前），并在整个核函数中使用该句柄。

4.4.3.2. 仅按引用传递组句柄

将组句柄传入函数时，建议按引用向函数传递组句柄。由于没有默认构造函数，组句柄必须在声明时初始化。不建议通过复制构造来创建组句柄。

4.4.4.1. 避免组创建风险

划分组是一项集合操作，组中的所有线程都必须参与。如果在并非所有线程都会到达的条件分支中创建组，可能导致死锁或数据损坏。

4.4.5.1. 同步

可以通过调用集合函数sync()来同步一个组。与__syncthreads()一样，sync()函数提供以下保证：

• 组中线程在同步点之前进行的所有内存访问（例如读取和写入），在同步点之后对组中所有线程都可见。

• 组中的所有线程都到达同步点之后，才允许任何线程越过该同步点继续执行。

以下示例展示了与__syncthreads()等效的cooperative_groups::sync()。

namespace cg = cooperative_groups;

cg::thread_block my_group = cg::this_thread_block();

// Synchronize threads in the block
cg::sync(my_group);

协作组可用于同步整个网格。从 CUDA 13 起，协作组不再能用于多设备同步。有关详细信息，请参阅大规模组一节。

有关同步的更多信息，请参阅协作组 API。

4.4.5.2. 屏障

协作组提供了一个类似于cuda::barrier的屏障 API，可用于更高级的同步。协作组屏障 API 与cuda::barrier在以下几个关键方面有所不同：

• 协作组屏障会自动初始化

• 组中的所有线程在每个阶段都必须到达屏障并在屏障处等待一次。

• barrier_arrive返回一个arrival_token对象，该对象必须传入对应的barrier_wait，并在其中被消耗，此后不能再次使用。

程序员在使用协作组屏障时必须注意避免以下风险：

• 在调用barrier_arrive之后、调用barrier_wait之前，组不能使用任何集合操作。

• barrier_wait仅保证组中的所有线程都已调用barrier_arrive。barrier_wait并不保证所有线程都已调用barrier_wait。

namespace cg = cooperative_groups;

cg::thread_block my_group = this_block();
cg::cluster_group cluster = this_cluster();

auto token = cluster.barrier_arrive();

// Optional: Do some local processing to hide the synchronization latency
     local_processing(my_group);

// Make sure all other blocks in the cluster are running and initialized shared data before accessing dsmem
cluster.barrier_wait(std::move(token));

4.4.6.1. 归约

reduce函数用于对指定组中每个线程提供的数据执行并行归约。必须提供下表所示运算符之一来指定归约类型。

归约会在硬件加速可用时使用硬件加速（要求计算能力为 8.0 或更高）。对于不具备硬件加速的较旧硬件，可以使用软件回退实现。只有 4 字节类型可以获得硬件加速。

有关归约的更多信息，请参阅协作组 API。

以下示例展示了如何使用cooperative_groups::reduce()在线程块范围内执行求和归约。

namespace cg = cooperative_groups;

cg::thread_block my_group = cg::this_thread_block();

int val = data[threadIdx.x];

int sum = cg::reduce(my_group, val, cg::plus<int>());

// Store the result from the reduction
if (my_group.thread_rank() == 0) {
   result[blockIdx.x] = sum;
}

4.4.6.2. 扫描

协作组包含可用于任意组大小的inclusive_scan和exclusive_scan实现。这些函数对指定组所包含的每个线程提供的数据执行扫描操作。

程序员可以选择指定上面的归约运算符表中列出的归约运算符。

namespace cg = cooperative_groups;

cg::thread_block my_group = cg::this_thread_block();

int val = data[my_group.thread_rank()];

int exclusive_sum = cg::exclusive_scan(my_group, val, cg::plus<int>());

result[my_group.thread_rank()] = exclusive_sum;

有关扫描的更多信息，请参阅协作组扫描 API。

4.4.6.3. 单线程调用

协作组提供了一个invoke_one函数，用于必须由单个线程代表一个组执行工作中的串行部分的情况。

• invoke_one从调用组中任意选择一个线程，并使用该线程以所提供的参数调用所提供的可调用函数。

• invoke_one_broadcast与invoke_one相同，只是调用结果还会广播到组中的所有线程。

不保证线程选择机制具有确定性。

以下示例展示了invoke_one的基本用法。

namespace cg = cooperative_groups;
cg::thread_block my_group = cg::this_thread_block();

// Ensure only one thread in the thread block prints the message
cg::invoke_one(my_group, []() {
   printf("Hello from one thread in the block!");
});

// Synchronize to make sure all threads wait until the message is printed
cg::sync(my_group);

可调用函数内部不允许在调用组内进行通信或同步。允许与调用组之外的线程进行通信。

4.4.7.1. 异步内存复制的对齐要求

只有当源为全局内存、目标为共享内存，且两者均至少按 4 字节对齐时，memcpy_async才是异步的。为获得最佳性能，建议共享内存和全局内存均按 16 字节对齐。

4.4.8.1. 何时使用cudaLaunchCooperativeKernel

cudaLaunchCooperativeKernel是一个 CUDA 运行时 API 函数，用于启动使用协作组的单设备核函数，专为执行需要线程块间同步的核函数而设计。此函数确保核函数中的所有线程都能在整个网格范围内同步并协作，而传统 CUDA 核函数只允许在各个线程块内部同步，无法做到这一点。cudaLaunchCooperativeKernel确保核函数启动是原子的，即如果 API 调用成功，则指定数量的线程块将在指定设备上启动。

良好的做法是首先通过查询设备属性cudaDevAttrCooperativeLaunch来确保设备支持协作启动：

int dev = 0;
int supportsCoopLaunch = 0;
cudaDeviceGetAttribute(&supportsCoopLaunch, cudaDevAttrCooperativeLaunch, dev);

如果设备 0 支持该属性，这会将supportsCoopLaunch设置为 1。仅支持计算能力为 6.0 及更高的设备。此外，还需要在以下环境之一中运行：

• 不使用 MPS 的 Linux 平台

• 使用 MPS 且设备计算能力为 7.0 或更高的 Linux 平台

• 最新的 Windows 平台