Как работает варп с атомной операцией?

Question

Нити в warp работают физически параллельно, поэтому, если один из них (называемый, поток X) запускает атомную операцию, что будет делать другое? Подождите? Означает ли это, что все потоки будут ждать, пока поток X будет помещен в атомную очередь, получит доступ (mutex) и сделает какой-то материал с памятью, который был защищен этим мьютеком и мьютексом realese после?

Есть ли способ принять другие потоки для некоторой работы, например, считывает некоторую память, поэтому атомная операция скроет ее латентность? Я имею в виду, 15 простаивающих потоков это .. не хорошо, я думаю. Атомный действительно медленный, не так ли? Как я могу ускорить его? Есть ли какой-нибудь шаблон для работы с ним?

Работает ли атомная операция с блокировкой общей памяти для банка или всей памяти? Например (без mutexs), есть __shared__ float smem[256];

• Резьба1 работает atomicAdd(smem, 1);
• Резьба2 работает atomicAdd(smem + 1, 1);

Те темы, работает с разными банками, но в целом общей памяти. Они управляют parralel или они будут поставлены в очередь? Есть ли какая-либо разница с этим примером, если Thread1 и Thread2 являются разделенными перекосами или общим?

Answer 1

Я считаю что-то вроде 10 вопросов. На это сложно ответить. Предлагается задать один вопрос на вопрос.

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

1. без условных (например, если ... то ... иначе) В этом случае все потоки выполняются та же команда, что происходит с атомарной инструкцией , Затем все 32 потока будут выполнять атомный, хотя и не обязательно в одном месте. Все эти атомы будут обрабатываться SM, и в некоторой степени будут сериализованы (они будут полностью сериализованы, если они будут обновлять одно и то же местоположение).
2. с условными обозначениями Например, предположим, что у нас было if (!threadIdx.x) AtomicAdd(*data, 1);. Тогда поток 0 выполнил бы атомный, а других - нет. Может показаться, что мы могли бы заставить остальных делать что-то еще, но выполнение блокировки стоп-кода не позволяет этого. Выполнение Warp сериализуется таким образом, что все потоки, принимающие путь if (true), будут выполняться вместе, и все потоки, выполняющие путь if (false), будут выполняться вместе, но истинные и ложные пути будут сериализованы. Так что опять-таки мы не можем иметь разные потоки в деформации, выполняющие разные инструкции одновременно.

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

Некоторые из ваших других вопросов, похоже, ожидают, что транзакции памяти будут завершены в конце цикла инструкций, в котором они возникли. Это не так. При использовании глобальной и совместно используемой памяти мы должны предпринять специальные действия в коде, чтобы гарантировать, что предыдущие транзакции записи видны другим потокам (которые можно утверждать как свидетельство того, что транзакция завершена.) Одним из типичных способов сделать это является использование барьерных инструкций, таких как __syncthreads() или __threadfence() Но без этих барьерных инструкций потоки не ждут завершения записи. A (операция, зависящая от a) read может останавливать поток. Обычно запись не может остановить поток.

Теперь давайте посмотрим на вопросы:

так что если один из них начать атомную операцию, что другие будут делать? Подождите?

Нет, они не ждут. Атомная операция отправляется в функциональный блок SM, который обрабатывает атомы, и все потоки продолжают вместе. Поскольку атом обычно подразумевает чтение, да, чтение может останавливать деформацию. Но потоки не ждут завершения атомной операции (т. Е. Записи). Однако последующее чтение этого местоположения могло бы задержать деформацию, снова, ожидая завершения атомной (записи). В случае глобального атома, который, как гарантируется, обновит глобальную память, он приведет к недействительности L1 в исходном SM (если включен) и L2, если они содержат это местоположение в качестве записи.

Есть ли способ взять другие потоки для некоторой работы, например, считывает некоторую память, поэтому атомная операция скроет ее латентность?

Не совсем по причинам, изложенным в начале.

Atomic действительно медленный, не так ли? Как я могу ускорить его? Есть ли какой-нибудь шаблон для работы с ним?

Да, атомарный может сделать программы намного медленнее, если они доминируют в деятельность (например, наивное сокращения или наивное гистограммировании.) Вообще говоря, способ ускорения атомарных операций, чтобы не использовать их, или использовать их экономно, таким образом, чтобы он не доминировал над программной деятельностью. Например, наивное сокращение будет использовать атом, чтобы добавить каждый элемент в глобальную сумму. Умное параллельное сокращение вообще не будет использовать атомы для работы, выполняемой в блоке threadblock. В конце сокращения потоковой блокировки один атом может использоваться для обновления частичной суммы threadblock в глобальную сумму. Это означает, что я могу выполнить быструю параллельную редукцию произвольно большого числа элементов, возможно, порядка 32 атомных добавок или меньше. Это скупость Атомикс будет в основном не будет заметно в общем исполнении программы, за исключением того, что она позволяет параллельное сокращение должно быть сделано в одном вызове ядра, а не 2.

Общая память: Работает ли они parralel или они будет поставлена ​​в очередь?

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

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