Как мне сделать обратное shfl.idx (т.е. разброс по варпу вместо сбора по деформации)?

Question

С помощью инструкции CUDA shfl.idx мы выполняем то, что по сути является сборкой внутри деформации: каждая линия обеспечивает базовую линию и исходную линию и получает данные исходной линии.

А как насчет обратной операции, разброс? Я имею в виду, не по памяти, а по дорожкам. Таким образом, каждая полоса обеспечивает базовую линию и полосу назначения, а для полос, на которые нацелен только ровно одна полоса, они получают значение целевой полосы; другие полосы заканчиваются неопределенным / произвольным значением.

Я почти уверен, что у PTX такого нет. Возможно, оно существует в SASS? Если нет, есть ли лучший способ реализовать это, чем, скажем, рассеяние на разделяемую память и загрузка из разделяемой памяти, как по индексу дорожек?

Answer 1

Все операции тасования определяются в терминах полосы, с которой нужно читать. Функции CUDA отображаются почти напрямую в инструкции ptx, которые сами отображаются почти напрямую в SASS. Все они являются вариациями операции "Сделать это значение доступным для чтения другими пользователями и чтения значения из заданной целевой полосы" с различными удобными способами указания целевой полосы.

В общем, вы должны попытаться перенастроить свою функцию, чтобы вам не понадобилась операция "разброс". Там нет инструкции, которая делает то, что вы хотите.

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

Answer 2

Как обстоят дела с сегодняшними графическими процессорами (Hopper и более ранние) — вы просто не знаете. Аппаратной поддержки межполосного рассеяния нет.

Итак, возможно, просто сделайте это простым способом - через общую память:

1. Получите рабочий буфер разделяемой памяти из 32 элементов
2. Пусть каждая дорожка записывает свои данные в буфер со смещением, равным предполагаемому месту назначения.
3. Синхронизировать (__syncthreads()Наверное)
4. Каждая дорожка считывает соответствующий буферный элемент

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