CUDA并行库Cooperative Groups

  1、Cooperative Groups

         在 CUDA 编程中，高效的并行算法往往需要线程协作(threads cooperate)以及共享数据(share data)来完成集体计算(collective computations)。要共享数据，线程间必然会涉及同步，而共享的粒度因算法而异，因此线程间的同步应尽量足够灵活，比如开发者可以显示地指定线程间同步，这样就可以确保程序的安全性、可维护性和模块化设计。

        Cooperative Groups（协同组）是CUDA 9.0引入的一个新概念，主要用于跨线程块（block）的同步。在CUDA 9.0之前，CUDA仅支持线程块内的同步，CUDA提供了2个原语操作：

__syncthreads()函数用于同步同一线程块内的所有线程

__syncwarp(unsigned mask=0xffffffff)函数用于同步线程束内的线程

        线程块级的同步并不能满足开发者的需求，在某些时候，开发者需要跨线程块同步，针对此问题，CUDA 9.0推出了Cooperative Groups机制，用于线程块内和跨线程块的同步。该机制为开发者提供了自定义线程组的方式，并提供了相应的同步函数，同时还包括一个新的kernel启动API（cudaLaunchCooperativeKernel），该API保证了Cooperative Groups同步的安全性。

       Cooperative Groups引入了一个新的数据结构：thread_block，即线程块。thread_block可以通过this_thread_block()进行获取并初始化。thread_block继承自更广义的线程组数据结构：thread_group 。thread_group 提供了如下函数：

void sync(); //同步组内的所有线程，这里X.sync()等价于__syncthreads()

unsigned size(); //获取组内的线程数目

unsigned thread_rank(); //获取线程的组内索引值（[0,size]）

bool is_valid(); //判断本组是否违背了任何APIconstraints（API限制）

        thread_block则提供如下特定线程块函数：

dim3 group_index(); //网格grid内3维索引（block索引）

dim3 thread_index(); //块block内3维索引（线程索引）

        相比__syncthreads()函数，使用X.sync()的好处在于避免了隐式同步隐患。

2、块间同步（网格同步）

        相比块内组，Cooperative Groups最强大的能力在于跨线程块同步，在CUDA 9.0之前，不同线程块仅能在kernel执行结束时同步，现在开发者可以通过grid_group 结构执行网格级同步。同步操作如下：

grid_group grid = this_grid();

grid.sync();

      不同于传统的<<<...>>>执行配置，网格级同步必须通过cudaLaunchCooperativeKernel API配置并启动kernel：

cudaError_t cudaLaunchCooperativeKernel( const T *func, //kernel函数指针 dim3 gridDim, dim3 blockDim, void **args, //kernel参数数组 size_t sharedMem = 0, cudaStream_t stream = 0 )

        特别注意为保证所有协同线程块能安全的常驻GPU，gridDim和blockDim的值需要慎重考虑，开发者可以通过计算SM的最大活跃线程块数目来最大化并行率：cudaOccupancyMaxActiveBlocksPerMultiprocessor( &numBlocksPerSm, my_kernel,numThreads , 0 ); // 初始化后运行Kernel

cudaLaunchCooperativeKernel((void*)my_kernel,deviceProp.multiProcessorCount*numBlocksPerSm, numThreads, args );

3、多GPU同步（多设备同步）

        类似网格级同步，多设备同步通过multi_grid_group 结构执行：

multi_grid_group multi_grid = this_multi_grid();

multi_grid.sync();

        并通过cudaLaunchCooperativeKernelMultiDevice API配置并启动kernel：

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

注意：

1. 该API将确保一个launch操作是原子的，例如当API调用成功时，相应数目的线程块在所有指定设备上launch成功。

2. 对于所有设备，该API调用的kernel函数必须是相同的。

3. 同一设备上的launchParamsList参数必须是相同的。

4. 所有设备的计算能力必须是相同的（major and minor versions）。

5. 对于所有设备，配置的网格大小（gridDim）、块大小（blockDim）和每个网格的共享内存大小必须是相同的。

6. 自定义的__device__,__constant__,__managed__全局变量在每个设备上都是独立实例化的，因此需要开发者对该类变量赋初值。

4、其他

        深入了解请参考官方文档：

Cooperative Groups: Flexible CUDA Thread Programming | NVIDIA Technical Blog