#cuda #logic #gpu
#cuda #opencl #графический процессор #gpgpu
Вопрос:
У меня вопрос о предикации ветвления в графических процессорах. Насколько я знаю, в графических процессорах они выполняют предикацию с помощью ветвей.
Например, у меня есть код, подобный этому:
if (C)
A
else
B
итак, если для завершения выполнения A требуется 40 циклов, а B — 50 циклов, если предположить, что для одного warp выполняются и A, и B, то требуется ли в общей сложности 90 циклов для завершения этой ветви? Или они перекрывают A и B, т. Е. когда выполняются некоторые инструкции A, затем ожидаются запросы к памяти, затем выполняются некоторые инструкции B, затем ожидаются запросы к памяти и так далее?
Спасибо
Комментарии:
1. Для тех, у кого возникнет соблазн отредактировать этот вопрос дальше, пожалуйста, обратите внимание, что предикация ветвления и предсказание ветвления — это не одно и то же……
2. я нашел хорошее описание здесь: yosefk.com/blog/simd-simt-smt-parallelism-in-nvidia-gpus.html
Ответ №1:
Все архитектуры с поддержкой CUDA, выпущенные на данный момент, работают как SIMD-машина. Когда в рамках основы происходит расхождение ветвей, оба пути кода выполняются всеми потоками в основе, при этом потоки, которые не следуют по активному пути, выполняют функциональный эквивалент NOP (кажется, я припоминаю, что к каждому потоку в основе прикреплен флаг условного выполнения, который позволяет маскировать неисполняющиеся потоки).
Итак, в вашем примере ответ «90 циклов», вероятно, является лучшим приближением к тому, что происходит на самом деле, чем альтернатива.
Комментарии:
1. Насколько я помню, в CUDA asm есть инструкция перехода, но переход будет приниматься всеми потоками warp.
2. Как ptx_isa.pdf: «Если потоки warp расходятся через зависящую от данных условную ветвь, warp последовательно выполняет каждый выбранный путь ветвления, отключая потоки, которые не находятся на этом пути, и когда все пути завершаются, потоки сходятся обратно к одному и тому же пути выполнения.». Итак, в PTX есть условная ветвь, но все потоки Warp должны принимать или не принимать эту ветвь одновременно, чтобы быть единообразными (для повышения производительности)
3. Скорее всего, это компромисс. Чем сложнее вы создаете предикацию ветвления и условное выполнение, тем больше транзисторов требуется для ее реализации. Графические процессоры предназначены для выполнения кода, который не имеет большого количества ветвей, и в этом сценарии имеет больше смысла использовать как можно больше транзисторов для частей GPU, которые выполняют вычисления, и меньше для частей, которые уменьшают задержку (планирование команд, кэш и т.д.).
4. Это неверно, условное выполнение обрабатывается за деформацию, а не за половину деформации. Также стоит отметить, что если условие ветвления не расходится в пределах деформации (например
if (threadIdx.x > 64), то расходящееся выполнение отсутствует.5. Транзакции с памятью обрабатываются в соответствии с warp также на архитектуре Fermi (текущая архитектура). Более старый GT200 (он же архитектура Tesla) обрабатывал транзакции с памятью за половину деформации. Половинные деформации нигде в Fermi не используются.