Что быстрее в CUDA: запись в глобальную память + __threadfence() или atomicExch() в глобальную память?

Question

Предполагая, что у нас есть много потоков, которые будут последовательно обращаться к глобальной памяти, какой вариант работает быстрее в целом? Я сомневаюсь, потому что __threadfence() учитывает все записи в общую и глобальную память, но записи объединяются. С другой стороны, atomicExch() учитывает только важные адреса памяти, но я не знаю, слиты записи или нет.

В коде:

array[threadIdx.x] = value;

Или же

atomicExch(&array[threadIdx.x] , value);

Благодарю.

Answer 1

На графические процессоры Kepler я бы поставил atomicExch так как атомы очень быстро на Кеплера. На Ферми это может быть стирка, но, учитывая, что у вас нет столкновений, atomicExch мог бы еще хорошо выступить.

Пожалуйста, проведите эксперимент и сообщите результаты.

Answer 2

Эти двое делают совершенно разные вещи.

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

__threadfence() задерживает текущий поток (и только текущий поток!), чтобы гарантировать, что любые последующие записи данным потоком действительно произойдут позже. В качестве таких, __threadfence() Сам по себе, без какого-либо дополнительного кода не очень интересно.

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

Обратите внимание, что ни один из них на самом деле не дает никаких гарантий относительно реального порядка выполнения потоков.