Как глобальные атомарные операции реализованы в Kepler? Я получил меньше производительности, используя gmem, а не атомные

Question

Я хотел бы знать реализацию глобальной атомики в Кеплере.

увидеть этот кусок кода:

1. if (threadIdx.x < workers) {
2.    temp = atomicAdd(dst, temp + rangeOffset);
3.    if (isLastPartialCalc(temp)) {                            
4.        atomicAdd(dst,-300000.0f);
5.    }
6. }

если я изменю строку 4 для этого:

* дст -= 300000.0f;

Производительность ниже! Изменение безопасно, так как больше потоков не будет писать на это значение (выходные данные одинаковы).

ядро с использованием atomic: ~883us ядро ​​с использованием gmem напрямую: ~903us

Я бегал несколько раз, и я всегда получаю ~20us штраф за изменение

ОБНОВЛЕНИЕ Похоже, что хранилище без использования атомарного всегда вызывает промах в L2, в то время как атомарная версия всегда производит попадание... поэтому я предполагаю, что попытка записи в какое-то место, которое было помечено (или что-то) с "атомарным", не разрешено в L2 и делает еще один запрос к gmem

Answer 1

Это обновление строки кэша, очевидно, более затратно (в вашем конкретном коде), чем просто второй глобальный атомарный доступ.

Единый глобальный атомарный доступ с одного SM к глобальной памяти на Kepler GK110 (например, K20) на самом деле довольно быстрый.

Как указано в официальном документе Кеплера, Кеплер улучшил скорость глобальной атомики по сравнению с Ферми.

Пропускная способность атомарной операции к общему адресу глобальной памяти повышается в 9 раз до одной операции за такт.

Answer 2

Атомика имеет семантику "забей и забудь". Это означает, что ядро ​​вызывает атомарную операцию и позволяет фактической атомарной операции выполняться кешем (не на SM), и ядро перейдет к следующей инструкции, не дожидаясь завершения фактической атомарной операции. Это работает только в том случае, если нет возвращаемого значения из атомарной операции, что имеет место в этом примере. Семантика "забей и забудь" позволяет СМ продолжать свои вычисления, разгружая вычисления атома в кэш.

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