__activemask() против __ballot_sync()

Question

После прочтения этого поста в блоге разработчиков CUDA я пытаюсь понять, когда безопасное \ правильное использование __activemask() на месте __ballot_sync(),

В разделе Active Mask Query авторы написали:

Это неверно, так как это приведет к частичной сумме вместо общей суммы.

и после, в разделе Оппортунистическое программирование на уровне деформации они используют функцию __activemask() так как:

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

Answer 1

Здесь нет __active_mask() в CUDA. Это опечатка (в статье блога). Так должно быть __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() будет исполнен в неотклоненном состоянии.

Однако в модели исполнения Вольта не только нет никаких гарантий, предоставляемых 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() на практике эти два выражения возвращают одно и то же значение.

В модели исполнения Вольта мы можем иметь случайное расхождение в statement_B(), Поэтому эти два выражения могут не возвращать один и тот же результат. Зачем?

__ballot_sync() Инструкция, как и все другие встроенные функции уровня CUDA 9+, имеющие параметр маски, имеют синхронизирующий эффект. Если у нас есть дивергенция с применением кода, если не может быть выполнен синхронизирующий "запрос", указанный аргументом mask (как в случае выше, когда мы запрашиваем полную конвергенцию), это будет представлять недопустимый код.

Однако, если мы имеем случайную дивергенцию (только для этого примера), __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); 

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