Подтвердить что ты не робот

CUDA: синхронизация потоков

Почти везде, где я читал о программировании с CUDA, есть упоминание о важности того, что все нити в warp делают то же самое.
В моем коде у меня есть ситуация, когда я не могу избежать определенного состояния. Это выглядит так:

// some math code, calculating d1, d2
if (d1 < 0.5)
{
    buffer[x1] += 1;  // buffer is in the global memory
}
if (d2 < 0.5)
{
    buffer[x2] += 1;
}
// some more math code.

Некоторые из потоков могут входить в один для условий, некоторые могут входить в оба, и другие могут не входить в них.

Теперь для того, чтобы все потоки снова вернулись к тому, чтобы "делать то же самое" после условий, следует ли их синхронизировать после условий с помощью __syncthreads()? Или это как-то происходит автоматически?
Могут ли две нити не делать то же самое, потому что одна из них - одна операция позади, таким образом разрушая ее для всех? Или есть некоторые за кулисами попытки заставить их сделать то же самое снова после ветки?

4b9b3361

ОТВЕТЫ

Ответ 1

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

Но разные перекосы не синхронизируются таким образом. Поэтому, если ваш алгоритм требует, чтобы определенные операции были завершены во многих искажениях, вам нужно будет использовать явные вызовы синхронизации (см. Руководство по программированию CUDA, раздел 5.4).

EDIT: реорганизовал следующие несколько абзацев, чтобы прояснить некоторые вещи.

Здесь есть две разные проблемы: синхронизация инструкций и видимость памяти.

  • __syncthreads() обеспечивает синхронизацию команд и обеспечивает видимость памяти, но только внутри блока, а не через блоки (CUDA Programming Guide, Приложение B.6). Он полезен для чтения-чтения в общей памяти, но не подходит для синхронизации доступа к глобальной памяти.

  • __threadfence() обеспечивает видимость глобальной памяти, но не выполняет никакой синхронизации команд, поэтому, по моему опыту, она имеет ограниченное использование (но см. пример кода в Приложении B.5).

  • Глобальная синхронизация команд невозможна в ядре. Если вам нужно f() выполнить все потоки перед вызовом g() в любом потоке, разделите f() и g() на два разных ядра и вызовите их поочередно с хоста.

  • Если вам просто нужно увеличивать общие или глобальные счетчики, рассмотрите возможность использования функции инкремента атома atomicInc() (Приложение B.10). В случае вашего кода выше, если x1 и x2 не являются глобально уникальными (для всех потоков в вашей сетке), неатомные приращения приведут к состоянию гонки, аналогичному последнему абзацу Приложения B. 2.4.

Наконец, имейте в виду, что любые операции с глобальной памятью и функции синхронизации, в частности (включая атоматику), являются плохими для производительности.

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

Ответ 2

Из раздела 6.1 Руководства по лучшей практике CUDA:

Любая команда управления потоком (если, switch, do, for, while) может существенно повлиять пропускную способность инструкции, заставляя нити одного и того же варпа расходиться; то есть, для выполнения различных путей выполнения. Если это произойдет, различные пути выполнения должен быть сериализован, увеличивая общее количество инструкций, выполненных для этого перекос. Когда все пути выполнения завершены, потоки сходятся назад к тому же пути выполнения.

Итак, вам не нужно ничего особенного делать.

Ответ 3

В ответе Габриэля:

"Глобальная синхронизация команд невозможна в ядре. Если вам нужно выполнить f() для всех потоков перед вызовом g() в любом потоке, разделите f() и g() на два разных ядра и вызовите их последовательно из хозяин".

Что, если причина, по которой вам нужны f() и g() в том же потоке, заключается в том, что вы используете память регистров, и хотите, чтобы вы регистрировали или делились данными из f, чтобы добраться до g? То есть, по моей проблеме, вся причина синхронизации между блоками заключается в том, что данные из f необходимы в g - и выходу из ядра потребуют большого количества дополнительной глобальной памяти для передачи данных регистра с f на g, что я избегать

Ответ 4

Ответ на ваш вопрос - нет. Вам не нужно ничего особенного делать. В любом случае, вы можете исправить это, вместо своего кода вы можете сделать что-то вроде этого:

buffer[x1] += (d1 < 0.5);
buffer[x2] += (d2 < 0.5);

Вы должны проверить, можете ли вы использовать общую память и доступ к глобальной памяти в объединенном шаблоне. Также убедитесь, что вы НЕ хотите писать в один и тот же индекс более чем в 1 поток.