Критические разделы CUDA, модель выполнения потоков / деформаций и решения компилятора NVCC

Question

Недавно я опубликовал этот вопрос о критическом разделе. Вот похожий вопрос. В этих вопросах данный ответ говорит, что это зависит от компилятора, если код "работает" или нет, потому что порядок различных путей выполнения зависит от компилятора.

Чтобы проработать остальную часть вопроса, мне нужны следующие выдержки из руководства по программированию CUDA:

1. ... Отдельные потоки, составляющие деформацию, начинаются вместе по одному и тому же адресу программы, но у них есть свой собственный счетчик адресов команд и состояние регистрации, и поэтому они могут свободно переходить и выполняться независимо....
2. Деформация выполняет одну общую инструкцию за раз, поэтому полная эффективность достигается, когда все 32 потока деформации согласовывают свой путь выполнения. Если потоки деформации расходятся через зависимую от данных условную ветвь, деформация последовательно выполняет каждый выбранный путь ветвления, отключая потоки, которые не находятся на этом пути, и когда все пути завершены, потоки сходятся обратно к одному и тому же пути выполнения...,
3. Контекст выполнения (счетчики программ, регистры и т. Д.) Для каждой деформации, обрабатываемой мультипроцессором, поддерживается на кристалле в течение всего срока службы деформации. Следовательно, переключение с одного контекста выполнения на другой не требует затрат, и при каждом времени выполнения команды планировщик деформации выбирает деформацию, в которой есть потоки, готовые выполнить свою следующую инструкцию (активные потоки деформации), и выдает инструкцию этим потокам.,

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

Теперь возникает вопрос: компилятор не должен всегда размещать ветви в порядке, написанном пользователем?, есть ли высокоуровневый способ обеспечить порядок? Я знаю, что компилятор может оптимизировать, переупорядочивать инструкции и т. Д., Но он не должен фундаментально изменять логику кода (да, есть исключения, например, доступ к памяти без ключевого слова volatile, но именно поэтому ключевое слово существует, предоставить контроль пользователю).

---

редактировать

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

Answer 1

Я считаю, что порядок выполнения не устанавливается и не гарантируется компилятором CUDA. Это аппаратное обеспечение, которое его устанавливает - насколько я помню.

Таким образом,

компилятор не должен всегда размещать ветки в порядке, написанном пользователем?

В любом случае он не контролирует порядок исполнения

Есть ли высокий уровень для обеспечения порядка?

Просто инструкции по синхронизации как __syncthreads(),

Компилятор... не должен принципиально менять логику кода

Семантика кода CUDA отличается от семантики кода C++. Последовательное выполнение ветвей if не является частью семантики.

Я понимаю, что этот ответ может быть неудовлетворительным для вас, но так обстоит дело, к лучшему или к худшему.