Могу ли я использовать __syncthreads() после сброса потоков?

Можно ли использовать __syncthreads() в блоке, где я целенаправленно удалял потоки с помощью return?

В документации указано, что __syncthreads() должен быть вызван каждым потоком в блоке, иначе он приведет к взаимоблокировке, но на практике я никогда не испытывал такого поведения.

Пример кода:

__global__ void kernel(float* data, size_t size) {
    // Drop excess threads if user put too many in kernel call.
    // After the return, there are `size` active threads.
    if (threadIdx.x >= size) {
        return;
    }

    // ... do some work ...

    __syncthreads(); // Is this safe?

    // For the rest of the kernel, we need to drop one excess thread
    // After the return, there are `size - 1` active threads
    if (threadIdx.x + 1 == size) {
        return;
    }

     // ... do more work ...

    __syncthreads(); // Is this safe?
}

Ответ 1

Ответ на короткий вопрос: "Нет". Отклонение ветки уровня Warp вокруг команды __syncthreads() вызовет тупик и приведет к зависанию ядра. Ваш пример кода не гарантированно будет безопасным или правильным. Правильный способ реализации кода будет таким:

__global__ void kernel(...)

    if (tidx < N) {
        // Code stanza #1
    }

    __syncthreads();


    if (tidx < N) {
        // Code stanza #2
    }

    // etc
}

так что команды __syncthreads() выполняются безоговорочно.

EDIT: просто добавьте немного дополнительной информации, подтверждающей это утверждение, вызовы __syncthreads() будут скомпилированы в инструкцию PTX bar.sync для всех архитектур. Руководство PTX2.0 (p133) сообщает bar.sync и включает следующее предупреждение:

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

Поэтому, несмотря на любые утверждения об обратном, небезопасно иметь условное разветвление вокруг вызова __syncthreads(), если вы не можете быть на 100% уверены, что каждый поток в любом заданном warp следует за тем же кодом путь и не может иметь место расхождение.

Ответ 2

Обновление Compute Capability 7.x(Volta):

С введением независимого потока Scheduling среди потоков в warp, CUDA, наконец, более строг на практике, теперь сопоставляя документированное поведение. Из Руководство по программированию:

Хотя __syncthreads() последовательно документируется как синхронизация всех потоков в блоке потоков, Pascal и предыдущие архитектуры могут только обеспечивать синхронизацию на уровне основы. В некоторых случаях это позволяло преграждать успех, не выполняясь нитями, если хотя бы какая-то нить в каждой деформации достигла барьера. Начиная с Volta, встроенные в CUDA команды __syncthreads() и команды PTX.sync(и их производные) применяются в каждом потоке и, следовательно, не будут достигнуты, пока не будут достигнуты все не выходящие потоки в блоке. Код, использующий предыдущее поведение, скорее всего, будет блокирован и должен быть изменен, чтобы гарантировать, что все не выходящие потоки достигнут барьера.

Ниже приведен предыдущий ответ, который чередовался с предвольтайским поведением.

Обновление. Этот ответ может не добавить ничего в дополнение к когтям (в зависимости от вашего понимания предмета, я полагаю), но рискуя быть слишком подробным, я представляю информацию что помогло мне понять это лучше. Кроме того, если вас не интересует, как вещи могут работать "под капотом" или что может быть возможно за пределами официальной документации, здесь ничего не видно. Все сказанное, я до сих пор не рекомендую делать предположения, выходящие за рамки официально зарегистрированных, особенно в среде, которая надеется поддерживать несколько или будущие архитектуры. Прежде всего я хотел бы отметить, что, хотя это явно называется плохой практикой Руководство по программированию CUDA, фактическое поведение __syncthreads() может быть несколько отличным от того, как это описано, и для меня это интересно. Последнее, что я хочу, это распространять дезинформацию, поэтому я открыт для обсуждения и пересмотра своего ответа!

Несколько важных фактов

Для этого ответа нет TL; DR, так как существует слишком много возможностей для неправильной интерпретации, но вот некоторые важные факты для начала:

__syncthreads() ведет себя как барьер для перекосов в блоке, а не всех нитей в блоке, хотя при использовании его в качестве рекомендации он составляет одно и то же.
Если какой-либо поток в warp выполняет инструкцию PTX bar (например, от _syncthreads), это похоже на то, что все нити в warp имеют.
Когда вызывается bar.sync (как генерируется instrinsic __syncthreads()), количество прибытий для этого блока и барьера увеличивается на размер основы. Вот как достигнуты предыдущие пункты.
Расхождение потоков (несколько путей) обрабатывается путем сериализации выполнения ветвей. Порядок сериализации - это фактор, который может вызвать проблемы.
Нити в пределах основы не синхронизируются с помощью __syncthreads(). Инструкция не заставит варп останавливаться и ждать потоков на расходящихся дорожках. Выполнение ветвей сериализуется, поэтому только когда ветки воссоединяются или код завершается, потоки в warp затем повторно синхронизируются. До этого ветки выполняются последовательно и независимо. Опять же, только один поток в каждом деформации блока должен нажать __syncthreads() для продолжения выполнения.

Эти утверждения поддерживаются официальной документацией и другими источниками.

Интерпретация и документация

Так как __syncthreads() действует как барьер для перекосов в блоке, а не всех потоков в блоке, как описано в Руководстве по программированию, кажется, что простой ранний выход был бы хорошим , если в по крайней мере одна нить в каждой деформации попадает в барьер. (Но это не означает, что вы не можете вызывать взаимоблокировки с внутренним!) Это также предполагает, что __syncthreads() всегда будет генерировать простую инструкцию bar.sync a; PTX и что семантика этого тоже не изменится, поэтому не делайте этого в производстве.

Одно интересное исследование, с которым я столкнулся, действительно исследует, что происходит, когда вы противоречите рекомендациям CUDA Programming Guide, и они обнаружили, что в то время как действительно возможно вызвать тупик, злоупотребляя __syncthreads() в условных блоках, не все использование внутреннего кода в условном коде будет делать это. Из раздела D.1 в документе:

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

Это утверждение согласуется с бит документацией по PTX, указанным talonmies. В частности:

Барьеры выполняются по принципу варпа, как если бы все нити в основе были активны. Таким образом, если какой-либо нить в warp выполняет инструкцию bar, это похоже на то, что все нити в warp выполнили инструкцию bar. Все потоки в деформации задерживаются до тех пор, пока барьер не завершится, а счетчик прибытия для барьера увеличивается на величину деформации (не количество активных нитей в деформации).

Из этого ясно, почему необязательный счетчик потоков b в инструкции bar.sync a{, b}; должен быть кратным размеру warp - всякий раз, когда один поток в warp выполняет команду bar количество прибытий увеличивается на размер основы, а не на количество нитей в основе, которые действительно попадают в барьер. Потоки, которые заканчиваются раньше (следуют по другому пути), фактически считались полученными в любом случае. Теперь следующее предложение в цитированном отрывке говорит, что не следует использовать __syncthreads() в условном коде, если только "известно, что все потоки одинаково оценивают условие (варп не расходится)". Это, по-видимому, слишком строгая рекомендация (для текущей архитектуры), предназначенная для обеспечения того, чтобы подсчет прибытий фактически отражал реальное количество потоков, попавших в барьер. Если хотя бы одна нить, поражающая барьер, увеличивает количество прибытий для всего варпа, у вас может быть немного больше гибкости.

Нет никакой двусмысленности в документации PTX, что команда bar.sync a;, сгенерированная __syncthreads(), ожидает, что все потоки в текущем массиве (блоке) совместного взаимодействия достигнут барьера a. Тем не менее, дело в том, что в настоящее время определяется "все потоки", увеличивая количество прибытий в кратном размере warp при каждом ударе барьера (по умолчанию, когда b не указывается). Эта часть не является undefined поведением, по крайней мере, не с версией ISA версии 4.2 Parallel Thread.

Имейте в виду, что в warp могут быть неактивные потоки даже без условного - "последние потоки блока, число потоков которых не кратно размеру основы". (Заметки архитектуры SIMT). Тем не менее __syncthreads() не запрещается в таких блоках.

Примеры

Ранняя версия выхода 1:

__global__ void kernel(...)

    if (tidx >= N)
        return;      // OK for <32 threads to hit this, but if ALL
                     // threads in a warp hit this, THEN you are deadlocked
                     // (assuming there are other warps that sync)

    __syncthreads(); // If at least one thread on this path reaches this, the 
                     // arrival count for this barrier is incremented by 
                     // the number of threads in a warp, NOT the number of 
                     // threads that reach this in the current warp.
}

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

Ранняя версия выхода 2:

__global__ void kernel(...)

    if (tidx < N) {
        // do stuff

        __syncthreads();
    }
    // else return;
}

По-прежнему нет взаимоблокировки, если у вас есть хотя бы один поток в warp, который попадает в барьер, но в этом случае важна последовательность выполнения ветвей? Я так не думаю, но, вероятно, это плохая идея, требующая определенного порядка выполнения.

Документ демонстрирует это в более сложном примере по сравнению с тривиальным ранним выходом, который также напоминает нам о том, чтобы быть осторожным в отношении расхождения в варпе. Здесь первая половина warp (идентификатор потока tid в [0,15]) записывается в некоторую разделяемую память и выполняет __syncthreads(), тогда как другая половина (поток id tid в [16,31]) также выполняет __syncthreads(), но теперь читается из разделов разделяемой памяти, написанных первой половиной основы. Сначала игнорируя тест на разделяемую память, вы можете ожидать тупика на любом барьере.

// incorrect code to demonstrate behavior of __syncthreads
if (tid < 16 ) {
  shared_array[tid] = tid;
  __syncthreads();
}
else {
  __syncthreads();
  output[tid] =
    shared_array[tid%16];
}

Нет взаимоблокировки, что означает, что __syncthreads() не синхронизирует расходящиеся потоки в пределах основы. Разновидные пути кода сериализуются в основе, и для прохождения вызова к __syncthreads() работает только один поток в кодовом пути.

Однако бит разделяемой памяти показывает, где может возникнуть какое-то непредсказуемое поведение. Вторая половина warp не получает обновленные значения из первой половины, потому что ветвь расходимости сериализовала выполнение warp и блок else был выполнен первым. Таким образом, функция работает неправильно, но также показывает, что __syncthreads() не синхронизирует расходящиеся потоки в warp.

Резюме

__syncthreads() не ждет всех нитей в деформации, и приход одного потока в warp эффективно подсчитывает всю деформацию как достижение барьера. (Современная архитектура).

В условном коде может быть опасно использовать __syncthreads() из-за сериализации расходящегося потока.

Использовать внутреннее выражение в условном коде, только если вы понимаете, как оно работает, и как обрабатывается разветвление (которое происходит внутри warp).

Обратите внимание, что я не сказал, чтобы идти вперед и использовать __syncthreads() таким образом, чтобы он не был документирован.