В CUDA нет __active_mask(). Это опечатка (в статье блога). Это должно быть __activemask().
__activemask() - это только запрос . Он задает вопрос "какие потоки в варпе в настоящее время выполняют эту инструкцию в этом цикле?" что равносильно тому, чтобы спросить «какие потоки в варпе в данный момент сходятся?»
Не влияет на конвергенцию. Это не приведет к схождению потоков. У него нет режима синхронизации варпа.
__ballot_sync(), с другой стороны, имеет сходящееся поведение (согласно предоставленному mask).
Основное различие здесь должно рассматриваться в свете модели исполнения деформации Вольта. Volta и другие из-за аппаратных изменений в механизме исполнения деформации могут поддерживать потоки в деформации, расходящиеся в большем числе сценариев и в течение более длительного времени, чем в предыдущих архитектурах.
Дивергенция, на которую мы ссылаемся здесь, является случайной дивергенцией из-за предыдущего условного исполнения. Принудительное расхождение из-за явного кодирования идентично до или после Вольта.
Давайте рассмотрим пример:
if (threadIdx.x < 1){
statement_A();}
statement_B();
Предполагая, что размер нити блока X больше 1, statement_A() находится в области принудительного расхождения. Деформация будет в дивергентном состоянии при выполнении statement_A().
А как насчет statement_B()? Модель выполнения CUDA не делает никаких конкретных заявлений о том, будет ли деформация в дивергентном состоянии или нет при выполнении statement_B(). В среде исполнения, предшествовавшей Volta, программисты, как правило, ожидали, что в закрывающей фигурной скобке предыдущего оператора if есть какая-то перекос конверсии (хотя CUDA не гарантирует этого). Поэтому общее ожидание состоит в том, что statement_B() будет выполнено в неотклоненном состоянии.
Однако в модели исполнения Volta не только нет никаких гарантий, предоставляемых CUDA, но на практике мы можем наблюдать деформацию в отклоненном состоянии в statement_B(). Дивергенция в statement_B() не требуется для правильности кода (тогда как это требуется в statement_A()), а также конвергенция в statement_B() не требуется моделью исполнения CUDA. Если в statement_B() есть расхождение, которое может возникнуть в модели исполнения Volta, я имею в виду это как случайное расхождение. Это расхождение возникает не из-за какого-то требования кода, а из-за какого-то предыдущего поведения условного выполнения.
Если у нас нет расхождения в statement_B(), то эти два выражения (если они были в statement_B()) должны возвращать один и тот же результат:
int mask = __activemask();
и
int mask = __ballot_sync(0xFFFFFFFF, 1);
Таким образом, в случае перед напряжением, когда мы обычно не ожидаем расхождения на statement_B(), на практике эти два выражения возвращают одно и то же значение.
В модели исполнения Volta у нас может быть случайная дивергенция в statement_B(). Поэтому эти два выражения могут не возвращать один и тот же результат. Почему?
Инструкция __ballot_sync(), как и все другие встроенные функции уровня CUDA 9+, имеющие параметр маски, имеет синхронизирующий эффект . Если у нас есть принудительное расхождение с кодом, если не удается выполнить синхронизирующий «запрос», указанный аргументом маски (как в случае выше, когда мы запрашиваем полную конвергенцию), это будет представлять недопустимый код.
Однако, если у нас случайная дивергенция (только для этого примера), семантика __ballot_sync() должна сначала реконвертировать деформацию по крайней мере в той степени, в которой аргумент маски запрашивает , затем выполните запрошенную операцию голосования.
Операция __activemask() не имеет такого поведения реконвергенции. Он просто сообщает о потоках, которые в настоящее время сходятся. Если некоторые потоки расходятся по по любой причине , они не будут указаны в возвращаемом значении.
Если затем вы создали код, выполняющий некоторую операцию на уровне деформации (например, уменьшение суммы на уровне деформации, как предлагается в статье блога), и выбрали потоки для участия на основе __activemask() против __ballot_sync(0xFFFFFFFF, 1), вы могли бы возможно получить другой результат, при наличии случайной дивергенции. Реализация __activemask(), при наличии случайной дивергенции, будет вычислять результат, который не включает все потоки (то есть она будет вычислять «частичную» сумму). С другой стороны, реализация __ballot_sync(0xFFFFFFFF, 1), поскольку она сначала устранила бы случайное расхождение, вынудила бы участвовать все потоки (вычисляя «общую» сумму).
Пример и описание, аналогичные приведенным здесь, приведены в листинге 10 в статье блога.
Пример правильного использования __activemask приведен в статье в блоге "Оппортунистическое программирование на уровне деформации", здесь:
int mask = __match_all_sync(__activemask(), ptr, &pred);
это утверждение говорит "скажите мне, какие потоки сходятся" (т.е. запрос __activemask()), а затем "используйте (как минимум) эти потоки для выполнения операции __match_all. Это совершенно законно и будет использовать все, что угодно в этот момент потоки сходятся. По мере продолжения этого примера листинга 9 mask, вычисленный на предыдущем шаге, используется в единственном другом примитиве, совместимом с варпом:
res = __shfl_sync(mask, res, leader);
(что происходит сразу после фрагмента условного кода). Это определяет, какие потоки доступны, а затем вынуждает использовать эти потоки независимо от того, какая случайная дивергенция могла существовать, для получения предсказуемого результата.