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

Question

Руководство по программированию CUDA утверждает, что

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

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

Я представляю, что могут быть всевозможные случаи, когда вам может понадобиться это сделать; например, если у вас есть бинарная маска и вам нужно условно применить определенную операцию к пикселям. Скажем, if (mask(x, y) != 0), затем выполните код, который включает __syncthreads(), иначе ничего не делайте. Как это будет сделано?

Answer 1

Если вам нужно пройти этот маршрут, вы можете разделить тело на две фазы:

if (condition)
{
    // code before sync
}
__syncthreads();
if (condition) // or remember a flag or whatever
{
    // code after sync
}

В качестве альтернативы вы можете использовать условие, чтобы установить флаг, который отключает определенные операции, например, если вы вычисляетедельта-обновление вы можете сделать следующее:

// *ALL* compute a delta update, those threads that would have failed the condition
// simply compute garbage.
// This can include syncthreads
if (condition)
    // apply update

Answer 2

Начиная с версии 3.0, вы можете использовать функции голосования деформации, чтобы выполнить то, что __syncthreads не может:

Функции голосования деформации поддерживаются только с помощью вычислительных устройств 1.2

int __all (предикат int); предикат для всех нитей основы и возвращает ненулевое значение тогда и только тогда, когда предикат оценивается как ненулевой для все они.

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

без знака int __ballot (предикат int); оценивает предикат для всех потоков деформация и возвращает целое число которого N-й бит устанавливается тогда и только тогда, когда предикат оценивается как ненулевой для N-нить основы. это функция поддерживается только устройствами вычислительных возможностей 2.x.

В противном случае есть также атомарные побитовые функции

atomicAnd, atomicOr, atomicXor

См. Раздел B.11 Руководства по программированию cuda