2. Cooperative Groups

2.1. 介绍

Cooperative Groups是一个MXMACA编程模型的扩展，目的是组织线程之间的交互。Cooperative Groups允许开发者表达交互线程的粒度，帮助他们进行更有效的并行分解。

2.2. 编程模型概念

Cooperative Groups编程模型描述MXMACA thread blocks内部以及相互之间的同步模式。它为应用提供定义他们自己的线程分组的方法，以及同步它们的接口。还提供新的强制执行特定限制的launch APIs来保证同步的功能。这些原语在MXMACA中启用新的协作并行模型，包括生产者-消费者并行性，机会并行性，以及整个Grid的整体同步。

Cooperative Groups编程模型由以下元素构成：

• 表示协作线程分组的数据类型；

• 获取MXMACA launch API定义的隐式分组的操作（比如 thread blocks）；

• 将现有分组划分为新分组的操作集合；

• 数据移动和操作的算法集合（比如 memcpy_async， reduce， scan）；

• 同步组内所有线程的操作；

• 检查组属性的操作；

• 暴露低级别、特定于组且通常由硬件加速的操作集合。

Cooperative Groups中的主要概念是对象命名其中一部分的线程集。这种将组表示为一级程序对象的方式改进了软件组合，因为集合函数可以接收表示参与线程组的显式对象。该对象还使程序员的意图明确，从而消除了从而消除了代码脆弱的不合理架构假设、对编译器优化的不良限制以及与GPU的更好的兼容性。

要编写高效的代码，最好使用专门的group（通用会失去很多编译时优化），并通过引用将这些组对象传递给打算以某种写作方式使用这些线程的函数。

要使用Cooperative Groups，请包含头文件：

#include <maca_cooperative_groups.h>

并使用Cooperative Groups的命名空间：

using namespace cooperative_groups;
// 或者使用一个别名防止集体算法污染命名空间
namespace cg = cooperative_groups;

代码可以使用mxcc正常编译，但是如果你想使用memcpy_async， reduce或者scan的功能，并且你的host侧编译器的默认方言不是c++11或更高，你需要在命令行添加–std=c++11。

2.2.1. 组示例

为了说明组的概念，此示例尝试执行块范围的规约求和，如果没有cooperative groups，写代码的时候在实现上有隐藏的约束：

__device__ int sum(int *x, int n) {
    __syncthreads();
    return total;
}

__global__ void parallel_kernel(float *x) {
    // ...
    // 整个线程块都必须调用sum
    sum(x, n);
}

所有线程块的线程都必须到达__syncthreads() barrier，但是这个约束对于想使用sum的开发者来说是隐藏的。使用Cooperative Groups，一个更好的实现方式：

__device__ int sum(const thread_block& g, int *x, int n) {
    // ...
    g.sync();
    return total;
}

__global__void parallel_kernel(...) {
    // ...
    // 整个线程块都必须调用sum
    thread_block tb = this_thread_block();
    sum(tb, x, n);
}

2.3. 组类型

2.3.1. 隐式组

隐式组代表内核的启动配置。无论内核是如何编写的，它始终具有一定数量的thread、block和block维度、单个grid和grid维度。另外，如果使用多设备协同启动API，可以有多个grid（每个设备一个grid）。这些组为分解成更细粒度的组提供了一个起点，这些组通常是硬件加速的，并且更专门用于开发人员正在解决的问题。

尽管您可以在代码的任何位置创建隐式组，但这样做很危险。为隐式组创建句柄是一个集体操作–组中的所有线程都必须参与。如果组不是在所有线程都到达的分支中创建的，则可能导致死锁或者数据损坏。因此，建议您预先为隐式组创建一个句柄（尽可能早，在任何分支发生之前）并在整个内核中使用这个句柄。

2.3.1.1. Thread Block Group

⼤家对于这样的thread block定义的thread group已经⾮常了解了。Cooperative Group扩展引⼊⼀个新的数据类型， thread_block，来显示地表示kernel中的这一概念。

class thread_block;
// 如下⽅式构造：
thread_block g = this_thread_block();
// Public 成员函数：
static void sync(); // 同步组内的线程
static unsigned int thread_rank(); // 在[0, num_threads)区间内为调⽤线程排序
static dim3 group_index(); // 在启动的grid中当前block的3-D index
static dim3 thread_index(); // 在启动的block中当前thread的3-D index
static dim3 dim_threads(); // 以units为单位的启动的block的尺⼨
static unsigned int num_threads(); // 当前group中总的线程数
// 传统成员函数（别名）：
static unsigned int size(); // 当前group中总的线程数(num_threads的别名)
static dim3 group_dim(); // 启动的block的尺⼨(dim_threads的别名)

例子：

/// 从global内存加载数据到shared内存
__global__ void kernel(int *globalInput) {
    __shared__ int x;
    thread_block g = this_thread_block();
    // 选⼀个leader
    if (g.thread_rank() == 0) {
        // ⼀次从global中加载所有thread需要的数据到shared
        x = (*globalInput);
    }
    // 在加载数据到shared内存完毕之后，如果block中所有的线程需要使⽤这个数据，你需
    g.sync(); // 等价于 __syncthreads();
}
// 注意： group中所有的线程都必须参与集体操作，否则⾏为是未定义的。
// thread_block数据类型是从更加generic的类型thread_group派⽣出来的，thread_group
// 可以用来表⽰更⼴泛的group

2.3.1.2. Grid Group

这个group对象表示在⼀个单独的grid中启动的所有threads。除了sync()之外的接口任何时间都可以调用，但是为了整个grid同步，你需要使用cooperative launch API（mcLaunchCooperativeKernel）。

class grid_group;
// 如下构造⽅式：
grid_group g = this_grid();
// Public 成员函数：
bool is_valid() const; // 返回当前grid_group是否可以同步
void sync() const; // 同步当前group中所有的线程
static unsigned long long thread_rank(); // 调⽤线程在[0, num_threads)上的排序
static unsigned long long block_rank()； // 调⽤block在[0, num_blocks)上的排序
static unsigned long long num_threads(); // group中threads的总数⽬
static unsigned long long num_blocks(); // group中block的总数⽬
static dim3 dim_blocks(); // 以units为单位的启动的block的⼤⼩
static dim3 block_index(); // 在启动的grid中当前block的3-D index
// 传统成员函数（别名）
static unsigned long long size(); // 当前group中总的线程数(num_threads的别名)
static dim3 group_dim(); // 启动的grid的尺⼨(dim_blocks的别名)

2.3.1.3. Multi Grid Group

这个group对象表示⼀个multi-device cooperative launch⾥⾯所有devices的所有threads。与grid_group不同，所有的API需要你使⽤合适的启动API（mcLaunchCooperativeKernelMultiDevice）。

class multi_grid_group;
// 如下构造⽅式
// 内核必须通过cooperative multi-device API启动
multi_grid_group g = this_multi_grid();
// Public 成员函数：
bool is_valid() const; // 返回当前multi_grid_group是否可以同步
void sync() const; // 同步当前group中所有的线程
unsigned long long num_threads(); // group中threads的总数⽬
unsigned long long thread_rank(); // 调⽤线程在[0, num_threads)上的排序
unsigned int grid_rank(); // 调⽤线程在[0, num_grids)上的排序
// 传统成员函数（别名）
static unsigned long long size(); // 当前group中总的线程数(num_threads的别名)
// 描述: 所有devices不推荐使⽤multi_grid_group。

2.3.2. 显式组

2.3.2.1. Thread Block Tile

⼀个模板化版本的tiled group, 模板参数⽤来指定tile的size，这样在编译时知道size，才有可能实现更优化的执⾏。

template <unsigned int Size, typename ParentT = void>
class thread_block_tile;
// 如下构造⽅式
template <unsigned int Size, typename ParentT>
_CG_QUALIFIER thread_block_tile<Size, ParentT> tiled_partition(const ParentT &g)

Size必须是2的幂，并且不⼤于64。

Parent T 是分割出这个⼦group的⽗类型。它是⾃动推断的，但是存储这个信息的void值将会在group句柄中存储，⽽不是这个类型中。

// Public 成员函数
void sync() const; // 同步group内的线程
unsigned long long num_threads() const; // group中线程总数
unsigned long long thread_rank() const; // 调⽤线程在[0, num_threads)上的排序
unsigned long long meta_group_size() const; // 返回parent group 分割创建的组的数量
unsigned long long meta_group_rank() const; // 从⽗group划分的tiles集合中的组的线性秩
// （由meta_group_size限定）
T shfl_up(T var, int delta) const; // 参考Warp Shuffle Functions
T shfl_down(T var, int delta) const; // 参考Warp Shuffle Functions
T shfl_xor(T var, int delta) const; // 参考Warp Shuffle Functions
T any(int predicate) const; // 参考Warp Vote Functions
T all(int predicate) const; //参考Warp Vote Functions
T ballot(int predicate) const; // 参考Warp Vote Functions
T match_any(T val) const; // 参考Warp Match Functions
T match_all(T val, int &pred) const; // 参考Warp Match Functions
// 传统成员函数（别名）
unsigned long long size() const; // 当前group中总的线程数(num_threads的别名)

在C++11及更高版本标准机型编译时，shfl, shfl_up, shfl_down与shfl_xor这些函数接收任意类型的对象。意味着只要满⾜以下条件，就可以shuffle⾮整数型的类型：

• 具有可复制性：⽐如， is_trivially_copyable<T>::value == true

• sizeof(T) <= 32

例⼦：

// 下⾯的代码将会创建两组tiled groups, size分别是32 和 4：
// 后者将源代码编码在类型中，⽽前者将其存储在句柄中
thread_block block = this_thread_block();
thread_block_tile<32> tile32 = tiled_partition<32>(block);
thread_block_tile<4, thread_block> tile4 = tiled_partition<4>(block);

备注

这里要使用模板化的数据结构 thread_block_tile ，分组的size通过模板参数传递给 tiled_partition 而非函数参数。

Warp-Synchronous代码模型

开发者可能会写出之前隐式地假定warp size的warp-synchronize代码，并且根据这个数字编写代码。现在需要显式地指定这个数字。

__global__ void cooperative_kernel(...) {
    // 获取默认的“当前thread block” 分组
    thread_block my_block = this_thread_block();

    // 细分为32个thread的分组，tiled子分组
    // tiled子分组将父分组均匀地划分为相邻的thread集合
    // ——在这种情况下，每个分组有一个warp size数量的thread
    auto my_tile = tiled_partition<64>(my_block);

    // 下面的操作将会只被每个block的前64个thread组成的
    // tiled分组执行
    if（my_tile.meta_group_rank() == 0) {
        // ...
        my_tile.sync();
    }
}

Single thread分组

表示当前thread的分组可以通过 this_thread 这个函数获取：

thread_block_tile<1> this_thread();

下面的 memcpy_async API使用一个 thread_group，来从源到目标拷贝一个元素：

#include <maca_cooperative_groups.h>
#include <maca_async.h>

cooperative_groups::memcpy_async(cooperative_groups::this_thread(),
dest, src, sizeof(int));

2.3.2.2. Coalesced Groups

在MXMACA的SIMT架构下，在硬件层面，多处理器以一组64线程（称作warps）的方式执行。如果在应用代码中存在数据依赖的条件分支，导致一个warp内的thread岔道，那么warp串行地执行每个分支，同时disable掉不在当前执行分支上的线程。在分支上的，保持active的线程被称作coalesced。Cooperative Groups有能力去发现，并且创建一个包含所有的coalesced线程的分组。

通过 coalesced_threads() 构造这个分组的句柄是机会主义的。它返回那个调用点当时active thread组成的集合，并且不保证哪一个thread是被返回的（只要他们是active的），或者说它们在整个执行过程中将会保持coalesced（它们在执行过程中作为一个集体，但是之后也可能再次分岔）。

class coalesced_group;
// 通过如下⽅式构造：
coalesced_group active = coalesced_threads();
// Public 成员函数：
void sync() const; // 同步group内的线程
unsigned long long num_threads() const; // group中线程总数
unsigned long long thread_rank() const; // 调⽤线程在[0, num_threads)上的排序
unsigned long long meta_group_size() const; // 返回parent group 分割创建的组的数量
// 如果这个group是通过查询活动线程集创建的，
// ⽐如coalesced_threads(), meta_group_size()值为1
unsigned long long meta_group_rank() const; // 从⽗group划分的tiles集合中的
// （由meta_group_size限定）
T shfl_up(T var, int delta) const; // 参考Warp Shuffle Functions
T shfl_down(T var, int delta) const; // 参考Warp Shuffle Functions
T shfl_xor(T var, int delta) const; // 参考Warp Shuffle Functions
T any(int predicate) const; // 参考Warp Vote Functions
T all(int predicate) const; //参考Warp Vote Functions
T ballot(int predicate) const; // 参考Warp Vote Functions
T match_any(T val) const; // 参考Warp Match Functions
T match_all(T val, int &pred) const; // 参考Warp Match Functions
// 传统成员函数（别名）
unsigned long long size() const; // 当前group中总的线程数(num_threads的别名

在C++11及更高版本标准机型编译时，shfl, shfl_up, shfl_down与shfl_xor这些函数接收任意类型的对象。意味着只要满⾜以下条件，就可以shuffle⾮整数型的类型：

• 具有可复制性：⽐如，is_trivially_copyable<T>::value == true

• sizeof(T) <= 32

例子：

__global__ void kernel(int *gIn) {
    if (threadIdx.x == *gIn) {
        coalesced_group active = coalesced_threads();
        active.sync();
    }
}

Discovery Pattern

开发者通常需要与当前active的线程集合共事。没有对存在的线程进行假设，而是开发人员使用碰巧存在的线程。可以通过下面的“聚合warp中线程之间的原子增量”的例子看出来：

{
    unsigned int writemask = __activemask();
    unsigned int total = __popc(writemask);
    unsigned int prefix = __popc(writemask & __lanemask_lt());
    // 发现laneID最小的active thread
    int elected_lane = __ffs(writemask) - 1;
    int base_offset = 0;
    if (prefix == 0) {
        base_offset = atomicAdd(p, total);
    }
    base_offset = __shfl_sync(writemask, base_offset, elected_lane);
    int thread_offset = prefix + base_offset;
    return thread_offset;
}

使用Cooperative Groups重新写这个功能如下：

{
    using cg = namespace cooperative_groups;
    cg::coalesced_group g = cg::coalesced_threads();
    int prev;
    if (g.thread_rank() == 0) {
        prev = atomicAdd(p, g.num_threads());
    }
    prev = g.thread_rank() + g.shfl(prev, 0);
    return prev;
}

2.4. 组分割

2.4.1. tiled_partition

template <unsigned int Size, typename ParentT>
thread_block_tile<Size, ParentT> tiled_partition(const ParentT& g);

thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);

tiled_partition 方法是一组分割父组到一个一维的，行优先的，子分组的平铺。将一共创建 (size(parent) / tilesz) 个子分组，因此父组的size必须是可以被Size整除的数字。允许的父组的类型是 thread_block 或者 thread_block_tile。

实现可能导致调用的thread一直等到所有的父组的成员调用了这个操作，之后才恢复执行。功能限制于本地的硬件条件，1/2/4/8/16/32/64，并且 cooperative_groups::size(parent) 必须比参数Size大。

例子：

/// 下面的代码将会创建一个32-thread tile
thread_block block = this_thread_block();
thread_block_tile<32> tile32 = tiled_partition<32>(block);

我们也可以把这些组分割成更小的组，每个组的size是4：

auto tile4 = tiled_partition<4>(tile32);
// 或使用一个通用的组
// thread_group tile4 = tiled_partition(tile32, 4);

假设我们要包含下面的代码行：

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

那么这个表达将会被block中每第四个线程打印。

2.4.2. labeled_partition

coalesced_group labeled_partition(const coalesced_group& g, int label);
template <unsigned int Size>
coalesced_group labeled_partition(const thread_block_tile<Size>& g, int label);

labeled_partition 方法是一个分割父组到一个一维的子分组（组内线程都是coalesced）的操作集合。实现将会评估条件label然后把有同样label值的线程放到相同的组。

实现可能导致调用的thread一直等到所有的父组的成员调用了这个操作，之后才恢复执行。

备注

这个功能仍然在测试中以后可能会有轻微改动。

2.4.3. binary_partition

coalesced_group binary_partition(const coalesced_group& g, bool pred);
template<unsigned int Size>
coalesced_group binary_partition(const thread_block_tile<Size>& g, bool pred);

binary_partition() 方法是一个分割父组到一个一维的子分组（组内线程都是coalesced）的操作集合。实现将会评估条件pred然后把有同样pred值的线程放到相同的组。它是 labeled_partition() 的一种特殊形式，label只能为0或1。

实现可能导致调用的thread一直等到所有的父组的成员调用了这个操作，之后才恢复执行。

备注

这个功能仍然在测试中以后可能会有轻微改动。

2.5. 组操作集合

2.5.1. 同步

2.5.1.1. sync

cooperative_groups::sync(T& group);

sync 函数同步指定的分组内的线程。T可以是任何存在的分组类型，它们都支持同步操作。如果这个分组是一个grid_group或者一个multi_grid_group，内核必须使用合适的cooperative启动接口来启动。

2.5.2. 数据移动

2.5.2.1. memcpy_async

memcpy_async 是一个组级别的内存拷贝操作集合，它对非阻塞的全局到共享内存的内存交换应用硬件加速支持。对于分组内的特定的线程集合， memcpy_async 会通过单独的一个流水步骤移动指定数量的字节或者输入类型的元素。此外，为了获取最佳的性能表现，需要全局以及共享内存按照16字节对齐。只有当拷贝源是全局内存并且拷贝目标是共享内存，并且两块内存编码都必须是16，8或者4字节对齐的时候这个操作才是异步的，否则它就是一个通常意义上的内存拷贝操作。异步内存拷贝的数据应该只能在 wait 之后读取，表示把数据移动到共享内存的操作阶段已经结束了。

必须等待所有未处理的请求可能会失去一些灵活性（但是会获得简单性）。为了有效地重叠数据传输和执行，能够在等待和操作请求N的同时启动N+1 memcpy_async 请求很重要。

用法1

template <typename TyGroup, typename TyElem, typename TyShape>
void memcpy_async(
    const TyGroup &group,
    TyElem *__restrict__ _dst,
    const TyElem *__restrict__ _src,
    const TyShape &shape
);

执行shape字节数的数据拷贝。

用法2

template <typename TyGroup, typename TyElem,
typename TyDstLayout, typename TySrcLayout>
void memcpy_async(
    const TyGroup &group,
    TyElem *__restrict__ dst,
    const TyDstLayout &dstLayout,
    const TyElel *__restrict__ src,
    const TySrcLayout &srcLayout
);

执行 min(dstLayout, srcLayout) 个元素的拷贝。

代码生成要求： C++11

头文件 maca_async.h 需要被包含。

例子：

/// 这个例子在block上把elementsPerThreadBlock个有效数据从全局内存拷贝到
/// 限制大小的共享内存（elementsInShared）。
#include <maca_cooperative_groups.h>
#include <maca_async.h>

namespace cg = cooperative_groups;

__global__ void kernel(int *global_data) {
    cg::thread_block tb = cg::this_thread_block();
    const size_t elementsPerThreadBlock = 16 * 1024;
    const size_t elementsInShared = 128;
    __shared__ int local_smem[elementsInShared];

    size_t copy_count;
    size_t index = 0;
    while (index < elementsPerThreadBlock) {
        cg::memcpy_async(tb, local_smem, elementsInShared, global_data + index,
                        elementsPerThreadBlock - index);
        copy_count = min(elementsInShared, elementsPerThreadBlock - index);
        cg::wait(tb);
        // 操作local_smem
        index += copy_count;
    }
}

2.6. Grid 同步

在介绍Cooperative Groups之前，MXMACA编程模型只允许在内核完成边界处的thread block之间进行同步。内核边界带来了一种隐含的状态无效，以及潜在的性能影响。

举个例子，在具体的应用场景，应用有很多的小内核，每一个内核表示一个流水进程中的一个状态。对于当前的MXMACA编程模型这些内核的存在是必要的，以保证运行在一个确定的流水线状态的thread block在下一个流水线状态之前输出了下一个流水线状态可能需要的数据。在这样的场景下，提供全局thread block同步的能力将会允许应用重构以拥有当给定阶段完成时能在设备上同步的持久thread block。

想要一个内核的整个grid同步，你只需要简单的使用 grid.sync() 函数：

grid_group grid = this_grid();
grid.sync();

而且在启动内核的时候，使用MXMACA 运行时启动API mcLaunchCooperativeKernel 替代<<<…>>>执行配置语法是必要的。

例子：

要保证GPU上的thread block的共同驻留，启动的block的数量需要仔细思考。比如像下面的尽可能多的启动SM可以启动的thread block：

int device = 0;
mcDeviceProp deviceProp;
mcGetDeviceProperties(&deviceProp, dev);
mcLaunchCooperativeKernel((void*)my_kernel, deviceProp.multiProcessorCount,
numThreads, args);

或者，你可以通过使用占用计算器计算每个SM可以同时容纳多少block来最大限度地提高暴露的并行性，如下所示：

/// 下面的代码将会启动一个在默认stream上的最大限度占用GPU的内核
int numBlocksPerSm = 0；
// my_kernel启动的线程数量
int numThreads = 128;
mcDeviceProp deviceProp;
mcGetDeviceProperties(&deviceProp, dev);
mcOccupancyMaxActiveBlocksPerMultiProcessor(&numBlocksPerSm,
my_kernel, numThreads, 0);
// 启动
void* kernelArgs[] = { /* 添加内核参数 */ }；
dim3 dimBlock(numThreads, 1, 1);
dim3 dimGrid(deviceProp.multiProcessorCount * numBlocksPerSm, 1, 1);
mcLaunchCooperativeKernel((void*)my_kernel, dimGrid, dimBlock, kernelArgs);

推荐通过查询设备的 mcDevAttrCooperativeLaunch 属性来先确认设备支持cooperative launch：

int dev = 0;
int supportsCoopLaunch = 0;
mcDeviceGetAttribute(&supportsCoopLaunch, mcDevAttrCooperativeLaunch, dev);

如果设备0的属性显示支持cooperative launch， supportsCoopLaunch 将会被设为1。

2.7. 多设备同步

为了支持使用Cooperative Groups进行多设备之间的同步，必须使用MXMACA启动接口 mcLaunchCooperativeKernelMultiDevice。这个与现存的 MXMACA 接口明显不同的接口允许一个主机线程在多个设备上启动一个内核。除了 mcLaunchCooperativeKernel 的限制与保证之外，这个接口有其他的语法：

• 这个接口保证启动是原子的，比如：如果这个接口调用成功，提供的线程块的数量将会在所有指定的设备上启动。

• 通过这个接口启动的函数必须是相同的。驱动不会在这一点上有显示的检查，因为是不可行的。这一点需要应用来保证。

• 在提供的 mcLaunchParams 中没有映射到相同设备的两个条目。

• 这个启动的所有目标设备必须有相同的计算能力——主要版本以及次要版本。

• 所有设备上每个grid的block size, grid size以及共享内存的数量都必须相同。这意味着每个设备上可以启动的block数量的最大值是受有最少SM数量的设备的限制的。

• 任何用户定义的 __device__ ， __constant__ 或者 __managed__ 的出现在模块中的设备侧全局变量拥有正在启动的CU函数，并在每个设备上独立实例化。用户有责任给这些全局变量合适的初始化。

可以通过启用对等访问（通过对所有参与的设备设置 mcCtxEnablePeerAccess 或者 mcDeviceEnablePeerAccess ）。

启动参数应该使用一个结构体的数组定义（一个设备一个结构体对象），并且使用 mcLaunchCooperativeKernelMultiDevice 启动。

例子：

mcDeviceProp deviceProp;
mcGetDeviceCount(&numGpus);

// 每个设备的启动参数
mcLaunchParams *launchParams =
                (mcLaunchParams*)malloc(sizeof(mcLaunchParams) * numGpus);
mcStream_t *streams = (mcStream_t*)malloc(sizeof(mcStream_t) * numGpus);

// 启动过程中拷贝内核启动参数
// 也可以使用每个设备独立的内核参数拷贝，但是函数/内核的签名以及名字必须相同
void* kernelArgs = { /* 添加内核参数 */ }

for (int i = 0; i < numGpus; i++) {
    mcSetDevice(i);
    // 每个设备的stream，但是每个设备也可以使用默认的NULL stream
    mcStreamCreate(&streams[i]);
    // 循环其他设备以及设置 mcDeviceEnablePeerAccess 以获取一个更快的barrier的实现
}

// 所有设备必须相同的计算能力并且有相同的启动配置
// 这里查询设备0是必要的
mcGetDeviceProperties(&deviceProp[i], 0);
dim3 dimBlock(numThreads, 1, 1);
dim3 dimGrid(deviceProp.multiProcessorCount, 1, 1);
for (int i = 0; i < numGpus; i++) {
    launchParamsList[i].func = (void*)my_kernel;
    launchParamsList[i].gridDim = dimGrid;
    launchParamsList[i].blockDim = dimBlock;
    launchParamsList[i].sharedMem = 0;
    launchParamsList[i].stream = streams[i];
    launchParamsList[i].args = kernelArgs;
}
mcLaunchCooperativeKernelMultiDevice(launchParams, numGpus);

此外，就像grid级别的同步，设备侧代码看起来也很像：

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

但是代码需要用独立编译的方式，传递-rdc=true到mxcc。

推荐通过查询设备的 mcDevAttrCooperativeMultiDeviceLaunch 属性来先确认设备支持多设备cooperative launch：

int dev = 0;
int supportsMdCoopLaunch = 0;
mcDeviceGetAttribute(&supportsMdCoopLaunch,
mcDevAttrCooperativeMultiDeviceLaunch, dev);

如果设备0的属性显示支持多设备cooperative launch， supportsMdCoopLaunch 将会被设为1。