背景

最近看到一个代码cooperative_groups.this_grid().sync()很好奇，这里好好梳理一下

分析

以前block内部的同步是用syncthreads(), block之间没有提供同步的接口，这样是合理的，假如有block间同步API的话，如果block太多，block_n要等block_0算完退出后才能进入sm, 但是block_0为了同步又要等block_n,这样就锁死了，本质原因是因为gpu的逻辑和cpu不一样，gpu单个block寄存器的值不会暂存到显存里来切换block_0。那么问题来了，这个this_grid().sync() API咋用？

实验

#include <stdio.h>
#include <cuda_runtime.h>
#include <cooperative_groups.h>namespace cg = cooperative_groups;__global__ void kernel(int* data) {int tid = threadIdx.x + blockIdx.x * blockDim.x;data[tid] = tid * tid;cg::this_grid().sync();if (blockIdx.x == 0 && threadIdx.x == 0) {for (int i = 0; i < gridDim.x * blockDim.x; ++i) {printf("%d ", data[i]);}printf("\n");}}template<typename... Types>
inline void launch_coop(void(*f)(Types...),dim3 gridDim, dim3 blockDim, cudaStream_t stream,Types... args)
{void* va_args[sizeof...(args)] = { &args... };(cudaLaunchCooperativeKernel((const void*)f, gridDim, blockDim,va_args, 0, stream));
}int main() {const int N = 8;int* d_data;cudaMalloc((void**)&d_data, N * sizeof(int));dim3 block(8);dim3 grid(20000);launch_coop(kernel, grid, block, 0, d_data);cudaDeviceSynchronize();cudaError_t err = cudaGetLastError();printf("$$$$$$$$$$$$$$: %s \n",cudaGetErrorString(err));cudaFree(d_data);return 0;
}

结论

实验结果发现，如果想使用这个API，必须保证所有的block在加载初期就是可以全部加载到sm中的，如果不行这个kernel launch就会失败，报错“too many blocks in cooperative launch”， 这个就比较合理了。

思考

这个api主要是避免小应用需要启用多个kernel来同步数据，因为单个kernel的block间无法同步（这个说法不是很准确），启动多个kernel不仅耗时，而且来回读写也耗时，所以小任务如果需要同步，可以考虑用这个API， 但是这个玩意是不是很耗时就没验证过了，有兴趣的可以试试。