Как синхронизировать Cuda в разных ветках?
У меня есть данные, которые я хочу обработать с помощью Cuda C++, где я работаю с четырьмя пикселями за раз, все четыре пикселя имеют общий угол. Например, я бы работал с четырьмя + пикселями вместе для всех мест в поле:
------------
--------++--
--------++--
------------
------------
Поэтому, чтобы избежать коллизий памяти, я решил выполнить свою работу в четыре этапа, где каждый этап может выполняться параллельно, не сталкиваясь с другими потоками, поскольку никакие два потока не работают одновременно с одним и тем же пикселем:
// We break the threads up into four flights:
//
// 0: even X and even Y
// 1: odd X and even Y
// 2: even X and odd Y
// 3: odd X and odd Y
const int flight = ( threadIdx.x % 2 + ( ( threadIdx.y % 2 ) << 1 ) );
for (int idx = 0; idx < flight; ++idx) {
__syncthreads();
}
doWork( pixel[ threadIdx.x + threadIdx.y * blockDim.x ],
pixel[ threadIdx.x + 1 + threadIdx.y * blockDim.x ],
pixel[ threadIdx.x + 1 + (threadIdx.y + 1) * blockDim.x ],
pixel[ threadIdx.x + (threadIdx.y + 1) * blockDim.x ]);
for (int idx = 3; idx > flight; --idx) {
__syncthreads();
}
Цель состоит в том, чтобы разбить работу на четыре полета, которые синхронизируются друг с другом, примерно так:
Четный рейс X / Четный Y:
doWork(...);
__syncthreads("one");
__syncthreads("two");
__syncthreads("three");
Нечетный рейс X / четный Y:
__syncthreads("one");
doWork(...);
__syncthreads("two");
__syncthreads("three");
Четный рейс X / нечетный Y:
__syncthreads("one");
__syncthreads("two");
doWork(...);
__syncthreads("three");
Нечетный рейс X / нечетный Y:
__syncthreads("one");
__syncthreads("two");
__syncthreads("three");
doWork(...);
Однако я думаю, что __syncthreads() не будет делать то, что я хочу, поскольку мое приложение не работает должным образом.
Могу ли я как-то идентифицировать свои точки синхронизации, например по имени, чтобы в коде было три синхронизации, между парами полетов, как показано в моем примере выше, или есть лучший способ сделать это в Cuda?
Конфигурация запуска выглядит так:
const int32_t pixelBlockSize = <argument to function>;
const int32_t pixelGridSize = <argument to function>;
const size_t scratch = (pixelBlockSize * pixelBlockSize + 2) * sizeof( float );
const dim3 dimBlock( pixelBlockSize, pixelBlockSize );
const dim3 dimGrid( pixelGridSize, pixelGridSize );
CallKernel<<< dimGrid, dimBlock, scratch >>> ( ... )
Ответы 2
__syncthreads()является точка синхронизации. Невозможно синхронизировать несколько отдельных потоков через __syncthreads(). Каждый __syncthreads() является барьером, который заставляет поток каждый в блоке ждать, пока потоки все блока не достигнут точки __syncthreads(). Вы не можете иметь __syncthreads() в расходящихся ветвях. Все (не завершенные) потоки блока должен прибывают в каждый __syncthreads(); в противном случае поведение не определено. Хотя на уровне PTX были бы способы выполняет более точную синхронизацию барьеров, я не думаю, что это действительно ответ. Если я правильно понимаю вашу проблему, то все, что вы ищете, похоже,
doWork(…)
__syncthreads();
doWork(…)
__syncthreads();
doWork(…)
__syncthreads();
doWork(…)
Все потоки каждого блока запускают начальный doWork(…) параллельно. Вы ждете, пока все потоки не будут сделаны с этим. Затем вы запускаете следующий doWork(…) и так далее…
В общем, вы также можете взглянуть на библиотека кооперативных групп, который предлагает хороший уровень абстракции поверх базовых примитивов синхронизации CUDA.
@talonmies Я не понимаю, что ты пытаешься мне сказать.
Хотя ваш ответ был правильным пять лет назад, сейчас он вводит в заблуждение. Вся синхронизация деформации на уровне PTX, о которой вы упомянули, полностью представлена в диалекте CUDA C, а функциональные возможности совместных групп обеспечивают именно то, о чем спрашивает вопрос - произвольные рабочие группы потоков с независимым поведением синхронизации и коллективными операциями.
@talonmies На самом деле я не говорил о синхронизации на уровне деформации. ОП спрашивал о синхронизации на уровне блоков. По крайней мере, я так понял вопрос. Обратите внимание, что кооперативные группы — это, по сути, просто библиотечная оболочка для существующих примитивов CUDA C и PTX. Я не знаю способа определить произвольные группы потоков внутри блока с независимым поведением синхронизации, как вы предлагаете. По замыслу кооперативные внутриблочные группы ограничены границами деформации. Вы не можете, например, синхронизировать дорожку 31 варпа 0 с дорожкой 0 варпа 1, используя совместные группы…
Я не знаю ничего, что вы могли бы сделать с помощью кооперативных групп, чего вы не могли бы сделать и без них. И они имеют те же ограничения, что и простое решение CUDA. Просто с ними проще и менее подвержены ошибкам (что, конечно, хорошо). Однако, поскольку OP спрашивал о поведении, возможно, самого основного механизма синхронизации, который существует в CUDA, я не видел причин поднимать библиотеку совместных групп. Я согласен, что упомянуть о его существовании не помешает…
Согласно ответу Майкла Кензеля:
// We break the threads up into four flights:
//
// 0: even X and even Y
// 1: odd X and even Y
// 2: even X and odd Y
// 3: odd X and odd Y
const int flight = ( threadIdx.x % 2 + ( ( threadIdx.y % 2 ) << 1 ) );
auto Process = [&](const bool run) {
if ( run )
{
doWork( pixel[ threadIdx.x + threadIdx.y * blockDim.x ],
pixel[ threadIdx.x + 1 + threadIdx.y * blockDim.x ],
pixel[ threadIdx.x + 1 + (threadIdx.y + 1) * blockDim.x ],
pixel[ threadIdx.x + (threadIdx.y + 1) * blockDim.x ]);
}
};
Process( 0 == flight ); __syncthreads();
Process( 1 == flight ); __syncthreads();
Process( 2 == flight ); __syncthreads();
Process( 3 == flight );
Обратите внимание, что это может показаться довольно неэффективным способом ведения дел. По крайней мере, если я правильно понимаю, только четверть ваших потоков будет что-то делать в любой момент времени. Казалось бы, лучше вместо этого, чтобы каждый поток работал с другим набором пикселей в каждом «полете»?
@MichaelKenzel, да, но есть другой код (не показан), в котором задействованы все потоки.
Как выглядит ваша конфигурация запуска? Вы запускаете только один блок? Если вы запускаете несколько блоков, гарантируется ли, что наборы пикселей, к которым обращаются потоки из отдельных блоков, не пересекаются? Обратите внимание, что
__syncthreads()обеспечивает синхронизацию только на уровне блоков (подробности см. в моем ответе ниже). Если вам нужно убедиться, что все потоки во всем запуске ядра выполнены, прежде чем двигаться дальше, вы, скорее всего, захотите просто выполнить несколько запусков ядра, например, по одному на сборку…