Модель согласованности памяти CUDA 4.0 и глобальная память?

Обновление :Условие while()ниже оптимизируется компилятором, поэтому оба потока просто пропускают условие и входят в CS даже с флагом -O0. Кто-нибудь знает, почему компилятор это делает? Кстати, объявление глобальных переменных volatileприводит к зависанию программы по какой-то странной причине...

Я читал Руководство по программированию CUDA , но я все еще немного не понимаю, как CUDA обрабатывает непротиворечивость памяти по отношению к глобальной памяти. (Это отличается от иерархии памяти. )По сути, я выполняю тесты, пытаясь нарушить последовательную согласованность . Алгоритм, который я использую, это алгоритм Петерсона для взаимного исключения между двумя потоками внутри функции ядра:

flag[threadIdx.x] = 1; // both these are global
turn = 1-threadIdx.x;

while(flag[1-threadIdx.x] == 1 && turn == (1- threadIdx.x));
shared_gloabl_variable_x ++;

flag[threadIdx.x] = 0;

Это довольно просто. Каждый поток запрашивает критическую секцию, устанавливая свой флаг в единицу и проявляя любезность, передавая очередь другому потоку. При оценке while(), если другой поток не установил свой флаг, запрашивающий поток может безопасно войти в критическую секцию.Теперь тонкая проблема с этим подходом заключается в том, что если компилятор повторно -упорядочивает записи так, чтобы запись в turnвыполнялась до записи в flag. Если это произойдет, оба потока окажутся в CS одновременно. Это довольно легко доказать с помощью обычных Pthreads, поскольку большинство процессоров не реализуют последовательную согласованность. А как насчет графических процессоров ?

Оба этих потока будут находиться в одном и том же варпе. И они будут выполнять свои операторы в пошаговом режиме блокировки -. Но когда они достигают переменной turn, они записывают в ту же самую переменную, поэтому выполнение внутренней -деформации становится сериализованным (независимо от порядка ). Теперь в этот момент выигравший поток переходит к условию while или ждет, пока другой поток закончит свою запись, чтобы оба затем могли одновременно оценить while()? Пути снова разойдутся на while(), потому что только один из них выиграет, пока другой ждет.

После запуска кода он последовательно ломает SC. Значение, которое я прочитал, равно ВСЕГДА 1, что означает, что оба потока каким-то образом входят в CS каждый раз. Как это возможно (Графические процессоры выполняют инструкции в порядке )? (Примечание. :Я скомпилировал его с -O0, поэтому без оптимизации компилятора и, следовательно, без использованияvolatile).