cuda中的cooperative_groups

背景

最近看到一个代码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， 但是这个玩意是不是很耗时就没验证过了，有兴趣的可以试试。