синхронизация блока cuda

Question

У меня есть b количество блоков, и каждый блок имеет t количество потоков. Я могу использовать

 __syncthreads()

для синхронизации потоков, которые находятся в определенном блоке. например

__global__ void aFunction()
{
    for(i=0;i<10;i++)
    {
       //execute something
        __syncthreads();
    }
}

Но моя проблема - синхронизировать все потоки во всех блоках. Как я могу это сделать?

Answer 1

В CUDA 9 NVIDIA представляет концепцию кооперативных групп , позволяющую синхронизировать все потоки, принадлежащие этой группе.Такая группа может охватывать все потоки в сетке.Таким образом, вы сможете синхронизировать все потоки во всех блоках:

grid_group g = this_grid();
g.sync();

Вам необходим Pascal или более новая архитектура для синхронизации сеток.

Основные функции, такие как синхронизациягруппы, меньшие, чем блок потока, вплоть до степени детализации, поддерживаются на всех архитектурах, в то время как графические процессоры Pascal и Volta поддерживают новые группы синхронизации для всей сетки и для нескольких графических процессоров.

Источник: https://devblogs.nvidia.com/parallelforall/cuda-9-features-revealed/

---

До CUDA 9 не было собственного способа синхронизации всех потоков из всех блоков.Фактически, концепция блоков в CUDA заключается в том, что некоторые из них могут быть запущены только после того, как некоторые другие блоки уже закончили свою работу, например, если графический процессор, на котором он работает, слишком слаб, чтобы обрабатывать их все параллельно.

Если вы гарантируете, что не создаете слишком много блоков, вы можете попытаться синхронизировать все блоки между собой, например, путем активного ожидания с использованием атомарных операций.Это, однако, медленно, съедая ваш контроллер памяти GPU, считается «взломом» и его следует избегать.

Так что, если вы не ориентируетесь на архитектуру Pascal (или более новую), лучший способ, которым я могуПредлагаем просто завершить работу вашего ядра в точке синхронизации, а затем запустить новое ядро, которое продолжит вашу работу.В большинстве случаев он будет работать быстрее (или, по крайней мере, с одинаковыми скоростями), чем при использовании упомянутого хака.