核函数中不同block线程如何同步

我在项目中遇到一个问题。和写了一个核函数,在核函数中又调用了一个核函数。第二个核函数的形式如下:
gpu1<<<4,16>>>();
我发现在gpu1运行完后通常是一组block线程运行完了,但是其他线程很没有运行完,我该如何同步不同block的线程。在实际项目中核函数参数是dim3的2D数据。

The only methods provided by CUDA to synchronize threads in separate blocks is the kernel launch boundaries (beginning and ending of a kernel, all threads are synchronized at the beginning or ending of your kernel code) and CUDA cooperative groups.

https://devblogs.nvidia.com/cooperative-groups/