2012-06-25 5 views
4

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

Кто-нибудь знает, как это сильно повлияет на производительность. (Официальная документация не очень специфична). У кого-нибудь есть более эффективный способ обработки этой части?

+1

Это зависит от того, что каждая нить в warp обычно принимает разные пути из-за переключателя или у вас есть, например. один коммутатор, который намного более распространен, чем остальные. Шаблон определит, сколько разницы в деформации вы получаете. –

+1

Как я знаю, коммутаторы, как правило, сопоставляются с инструкцией короткого перехода не как последовательность ifs. В C++ коммутатор обычно более эффективен, чем несколько операторов if. Я думаю, вам нужно исследовать отображение команд CUDA в этом случае до окончательного предложения. – geek

+1

Что находится внутри каждого заявления 'case'? Это однолинейный? Или несколько строк кода? Проводка примера, вероятно, даст вам более конкретные ответы. – Pedro

ответ

1

Хороший способ избежать использования нескольких коммутаторов - это реализовать таблицу функций и выбрать функцию из таблицы по индексу, основанному на условии переключения. CUDA позволяет использовать указатели на функцию __device__ в ядрах.

+0

Спасибо, это похоже на то, что мне нужно! – gamerx

+0

Возможно, глупый вопрос, но какой эффект это повлияет на производительность? Это обойдется только в том, что нужно написать большой оператор switch-case, если я не пропущу точку ... – Bart

+6

Как это помогает уменьшить разницу в деформациях и связанные с этим штрафы за производительность? –

4

Графический процессор работает с потоками по группам 32, называемым перекосами. Всякий раз, когда различные потоки в warp проходят через разные пути в коде, GPU должен запускать весь warp несколько раз, один раз для каждого пути кода.

Для решения этой проблемы, называемой расхождением warp, вы хотите упорядочить свои потоки, чтобы потоки в заданной деформации проходили как можно больше различных путей кода. Когда вы это сделали, вам в значительной степени просто нужно укусить пулю и принять потерю производительности, вызванную любой оставшейся деформой. В некоторых случаях, возможно, не все, что вы можете сделать, чтобы упорядочить свои потоки. Если это так, и если разные кодовые пути являются важной частью вашего ядра или общей рабочей нагрузки, задача может оказаться непригодной для графического процессора.

Не имеет значения как вы реализуете различные пути кода. if-else, switch, предикация (в PTX или SASS), таблицы ветвей или что-то еще - если дело доходит до потоков в warp, работающих на разных путях, вы получаете удар по производительности.

Также не имеет значения, сколько потоков проходит по каждому пути, а всего общее количество различных путей в деформации.

Здесь another answer на этом, что идет немного подробнее.

Смежные вопросы