У меня есть b количество блоков, и каждый блок имеет t количество потоков. Я могу использовать
__syncthreads()
для синхронизации потоков, которые находятся в конкретном блоке. например
__global__ void aFunction()
{
for(i=0;i<10;i++)
{
//execute something
__syncthreads();
}
}
Но моя проблема заключается в синхронизации всех потоков во всех блоках. Как я могу это сделать?
В CUDA 9 NVIDIA представляет концепцию кооперативных групп, позволяющую синхронизировать все потоки, принадлежащие этой группе. Такая группа может охватывать все потоки в сетке. таким образом, вы сможете синхронизировать все потоки во всех блоках:
#include <cooperative_groups.h>
grid_group g = this_grid();
g.sync();
Вам нужен Pascal (возможность вычисления 60) или более новая архитектура для синхронизации сеток.
Базовые функциональные возможности, такие как группы синхронизации, меньшие, чем блок потока, до степени детализации, поддерживаются на всех архитектурах, в то время как графические процессоры Pascal и Volta обеспечивают новые группы синхронизации для всей сетки и для нескольких графических процессоров.
Источник: https://devblogs.nvidia.com/parallelforall/cuda-9-features-revealed/
До CUDA 9 не было собственного способа синхронизации всех потоков из всех блоков. Фактически, концепция блоков в CUDA заключается в том, что некоторые из них могут быть запущены только после того, как некоторые другие блоки уже закончили свою работу, например, если графический процессор, на котором он работает, слишком слаб, чтобы обрабатывать их все параллельно.
Если вы гарантируете, что не создаете слишком много блоков, вы можете попытаться синхронизировать все блоки между собой, например, активно ожидая с использованием атомарных операций. Это, однако, медленно, съедая ваш контроллер памяти GPU, считается "взломом" и его следует избегать.
Поэтому, если вы не ориентируетесь на архитектуру Pascal (или более новую), лучший способ, который я могу предложить, - просто завершить работу вашего ядра в точке синхронизации, а затем запустить новое ядро, которое продолжит вашу работу. В большинстве случаев он будет работать быстрее (или, по крайней мере, с одинаковыми скоростями), чем при использовании упомянутого хака.