Обновление Compute Capability 7.x (Volta):
С введением независимого планирования потоков среди потоков в деформации, CUDA наконец-то стала более строгой на практике, теперь она соответствует документированному поведению. Из Руководства по программированию :
Хотя __syncthreads () постоянно документируется как синхронизирующая все потоки в блоке потоков, Pascal и предыдущие архитектуры могли обеспечить синхронизацию только на уровне деформации. В некоторых случаях это позволяло преграде преуспеть, не будучи выполненным каждым потоком, пока по крайней мере некоторый поток в каждой деформации достиг барьера. Начиная с Volta, встроенные в CUDA __syncthreads () и инструкция PTX bar.sync (и их производные) применяются для каждого потока и, таким образом, не будут успешными, пока не будут достигнуты всеми не завершенными потоками в блоке. Код, использующий предыдущее поведение, вероятно, будет тупиковым и должен быть изменен, чтобы гарантировать, что все невыпущенные потоки достигнут барьера.
Ниже приведен предыдущий ответ, который рассказывал о поведении до-Вольта.
Обновление : Этот ответ не может ничего добавить к талонмам (полагаю, в зависимости от вашего понимания предмета), но рискуя оказаться слишком многословным, я представляю информацию, которая помог мне понять это лучше. Кроме того, если вас не интересует, как все может работать "под капотом" или что может быть возможно за пределами официальной документации, здесь нечего смотреть. Тем не менее, я до сих пор не рекомендую делать предположения сверх того, что официально задокументировано, особенно в среде, которая надеется поддерживать несколько или будущие архитектуры. Прежде всего, я хотел бы отметить, что, хотя это Руководство по программированию явно указывает на плохую практику, реальное поведение __syncthreads() может несколько отличаться от того, как оно описано, и для меня это интересно. , Последнее, что я хочу, это распространять дезинформацию, поэтому я открыт для обсуждения и пересмотра своего ответа!
Несколько важных фактов
Для этого ответа не существует TL; DR, поскольку существует слишком большой потенциал для неправильного толкования, но вот несколько важных фактов для начала:
__syncthreads() ведет себя как барьер для перекосов в блоке, а не для всех потоков в блоке, хотя при использовании в соответствии с рекомендациями это равносильно одному и тому же. - Если какой-либо поток в деформации выполняет инструкцию PTX
bar (например, из _syncthreads), это как если бы все имеют потоки в деформации.
- Когда вызывается
bar.sync (как генерируется инстинктом __syncthreads()), счетчик прибытия для этого блока и барьера увеличивается на размер деформации. Так набираются предыдущие очки.
- Расхождение потоков (несколько путей) обрабатывается путем сериализации выполнения ветвей. Порядок сериализации является фактором, который может вызвать проблемы.
- Потоки внутри основы не синхронизируются
__syncthreads(). Инструкция не приведет к остановке деформации и ожиданию потоков на расходящихся путях. Выполнение ветвлений сериализуется, поэтому только когда ветви присоединяются или код завершается, потоки в деформации затем повторно синхронизируются. До этого ветки работают последовательно и независимо. Опять же, только один поток в каждой деформации блока должен нажать __syncthreads() для продолжения выполнения.
Эти заявления подтверждаются официальной документацией и другими источниками.
Толкование и документация
Так как __syncthreads() действует как барьер для перекосов в блоке, а не во всех потоках в блоке, как это описано в Руководстве по программированию, кажется, что простой ранний выход будет в порядке , если при как минимум одна нить в каждой деформации попадает в барьер . (Но это не значит, что вы не можете вызывать взаимные блокировки с помощью встроенного!) Это также предполагает, что __syncthreads() всегда будет генерировать простую bar.sync a; команду PTX и что семантика этого также не изменится, поэтому не делайте этого в производстве.
Одно интересное исследование , с которым я столкнулся, на самом деле исследует, что происходит, когда вы идете вразрез с рекомендациями Руководства по программированию CUDA, и они обнаружили, что, хотя действительно возможно вызвать тупик, злоупотребляя __syncthreads() в условных блоках не все используют встроенное в условном коде. Из раздела D.1 в документе:
Руководство по программированию рекомендует использовать syncthreads () в условном коде, только если условие оценивается одинаково по всему блоку потока. Остальная часть этого раздела исследует поведение syncthreads (), когда эта рекомендация нарушается. Мы демонстрируем, что syncthreads () работает как барьер для перекосов, а не потоков. Мы показываем, что когда потоки деформации сериализуются из-за расхождения ветвей, любые syncthreads () на одном пути не ожидают потоков с другого пути, а только ожидают других деформаций, работающих в том же блоке потоков.
Это утверждение согласуется с битом документации PTX , цитируемой talonmies. В частности:
Барьеры выполняются на основе основы, как если бы все потоки в основе были активными. Таким образом, если какой-либо поток в деформации выполняет команду bar, то все потоки в деформации выполнили инструкцию bar. Все нити в деформации останавливаются до тех пор, пока барьер не завершится, и счет прибытия для барьера увеличивается на размер деформации (а не количество активных нитей в деформации).
Из этого ясно, почему необязательный счетчик потоков b в инструкции bar.sync a{, b}; должен быть кратным размеру деформации - всякий раз, когда один поток в деформации выполняет bar инструкцию счет прибытия увеличивается на размер основы, а не на количество нитей в основе, которые фактически достигают барьера . Потоки, которые заканчиваются рано (следуя по другому пути), эффективно учитывались как полученные Далее, в следующем предложении в цитируемом отрывке говорится, что не следует использовать __syncthreads() в условном коде, если только «не известно, что все потоки оценивают условие одинаково (деформация не расходится)». Похоже, что это слишком строгая рекомендация (для текущей архитектуры), предназначенная для того, чтобы гарантировать, что счетчик поступлений действительно отражает реальное количество потоков, попадающих в барьер. Если хотя бы одна нить, ударяющая о барьер, увеличивает количество поступлений для всей основы, у вас действительно может быть немного больше гибкости.
В документации по PTX нет двусмысленности, что инструкция bar.sync a;, сгенерированная __syncthreads(), ожидает, пока все потоки в текущем кооперативном массиве потоков (блоке) достигнут барьера a. Тем не менее, дело в том, что «все потоки» в настоящее время определяются путем увеличения количества поступлений в кратных размерах варпа при каждом попадании в барьер (по умолчанию, когда b не указано). Эта часть не является неопределенным поведением, по крайней мере, в случае параллельного выполнения потоков ISA версии 4.2.
Имейте в виду, что в деформации могут быть неактивные потоки даже без условного выражения - «последние потоки блока, число потоков которых не кратно размеру деформации». ( Замечания по архитектуре 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;
}
Все еще нет тупика, если хотя бы один поток в варпе поразил барьер,но важен ли порядок выполнения ветвей в этом случае?Я так не думаю, но, вероятно, плохая идея требовать определенного порядка выполнения. В документе это демонстрируется на более сложном примере по сравнению с тривиальным ранним выходом, который также напоминает нам о необходимости быть осторожными в отношении деформации.дивергенция.Здесь первая половина деформации (идентификатор потока tid в [0,15]) выполняет запись в некоторую разделяемую память и выполняет __syncthreads(), а другая половина (идентификатор потока 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() на уровне за деформацию требуется только один поток в пути кода. Однако общий ресурсбит памяти показывает, где какое-то непредсказуемое поведение может войти в это.Вторая половина деформации не получает обновленные значения из первой половины, поскольку дивергенция ветвей сериализовала выполнение деформации и блок else был выполнен первым .Таким образом, функция работает неправильно, но она также показывает, что __syncthreads() не синхронизирует расходящиеся потоки в деформации. Сводка
__syncthreads() не ожидает все потоки вдеформация, и появление единой нити в деформации фактически учитывает весь деформацию как достижение барьера.(Представленная архитектура).
Может быть опасно использовать __syncthreads() в условном коде из-за того, как сериализуется выполнение расходящихся потоков.
Используйте встроенное в условном коде, только если вы понимаете, как оноработает и как обрабатывается расхождение ветвей (которое происходит в пределах основы).
Обратите внимание, что я не говорил идти вперед и использовать __syncthreads() способом, несовместимым скак это задокументировано.