Могу ли я настроить синхронизацию потоков CUDA по сетке?

Question

Я использовал кооперативные группы для выполнения синхронизации между потоками внутри ядра на уровне устройства:

      __global__ void my_kernel(...){
   grid_group grid = this_grid();

   // ... 
   grid.sync();
   // ...
}

И это совместный запуск, который делает доступной эту синхронизацию:

      cudaLaunchCooperativeKernel((void *)my_kernel, grid_dime, block_dime, kernel_args, shmem_size, 0);

Вызов sync() синхронизирует все потоки в сетке. Мои тесты показывают, что для большого количества потоков в сетке этот подход требует очень много времени, особенно когда блоки размещаются более чем на одном SM.

Я думал об уменьшении накладных расходов на синхронизацию, выполняя ее выборочно. Есть ли способ настроить эту синхронизацию? Подобно синхронизации некоторых потоков в отличие от всех, например, синхронизация первого потока всех блоков потоков.