Модель памяти CUDA: зачем приобретать забор не требуется для предотвращения переупорядочения нагрузки нагрузку? ⇐ C++

[Anonymous]

Я читаю книгу «Программирование массовых параллельных процессоров» и заметил приведенные ниже фрагменты кода для достижения «Сканирования в стиле домино»: < /p>

Код: Выделить всё

if (threadIdx.x == 0) {
while(AtomicAdd(&flags[bid], 0) == 0) {}
// ???? why do I not need thread fence here (for acquire semantic) to prevent load-load reordering between the loads for flags[bid] and scan_value[bid]?
previous_sum = scan_value[bid];
scan_value[bid+1] = previous_sum + local_sum;
__threadfence(); // why the "release" fence here is sufficient?
atomicAdd(&flags[bid + 1], 1);
}

Согласно книге, между этими двумя записями требуется __threadfence () для обеспечения SCACN_VALUE [BID+1] записывается в глобальную память перед флагами [ Bid+1] увеличен, что, по -видимому, напоминает семантическую семантику в модели памяти C ++ (за исключением того, что __threadfence имеет семантику SEQ_CST).
Однако я не могу понять, почему Аналогичный __threadfence (или приобретать забор) не требуется между двумя чтениями, чтобы предотвратить переупорядочение нагрузки на загрузку (то есть scan_value [bid] , прочитав перед флагами [bid] )?

Подробнее здесь: https://stackoverflow.com/questions/794 ... d-reorderi