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

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

Для уточнения остальной части вопроса мне нужны следующие выдержки из Руководство по программированию CUDA :

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

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

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


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

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

1 Ответ

0 голосов
/ 09 января 2019

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

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

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

Все равно не контролирует порядок исполнения

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

Только инструкции по синхронизации , как __syncthreads().

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

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

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

...