cuda: накладные расходы на деформацию деформации против дополнительной арифметики

Question

Конечно, деформация деформации, через if а также switch Заявлений, следует избегать любой ценой на графических процессорах.

Но каковы издержки дивергенции деформации (планирование только некоторых потоков для выполнения определенных строк) по сравнению с дополнительной бесполезной арифметикой?

Рассмотрим следующий фиктивный пример:

версия 1:

__device__ int get_D (int A, int B, int C)
{
    //The value A is potentially different for every thread.

    int D = 0;

    if (A < 10)
        D = A*6;
    else if (A < 17)
        D = A*6 + B*2;
    else if (A < 26)
        D = A*6 + B*2 + C; 
    else 
        D = A*6 + B*2 + C*3;

    return D;
}

против

версия 2:

__device__ int get_D (int A, int B, int C)
{
    //The value A is potentially different for every thread.

    return  A*6 + (A >= 10)*(B*2) + (A < 26)*C + (A >= 26)*(C*3);
}

Мой реальный сценарий более сложный (больше условий), но это та же идея.

Вопросы:

Являются ли накладные расходы (в планировании) расхождения деформации настолько велики, что версия 1) медленнее, чем версия 2?

Для версии 2 требуется намного больше ALU, чем для версии 1, и большинство из них тратится на "умножение на 0" (только некоторые из условных выражений оцениваются как 1, а не как 0). Связывает ли это ценные ALU в бесполезных операциях, задерживая инструкции в других деформациях?

Answer 1

Конкретные ответы на подобные вопросы обычно трудно дать. Есть много факторов, которые влияют на анализ сравнения между двумя случаями:

• Вы говорите, что А потенциально различна для каждого потока, но степень, в которой это верно, будет фактически влиять на сравнение.
• В целом, ответ зависит от того, привязан ли ваш код к вычислениям или пропускной способности. (Если ваш код ограничен пропускной способностью, между этими двумя случаями может не быть разницы в производительности).
• Я знаю, что вы определили A, B, C, как целые числа, но, казалось бы, безобидные изменения, как их сделать float может значительно повлиять на ответ.

К счастью, есть инструменты профилирования, которые могут помочь дать четкие, конкретные ответы (или, возможно, указать, что между этими двумя случаями нет большой разницы). Вы проделали довольно хорошую работу по выявлению двух конкретных случаев, которые вас волнуют. Почему бы не сравнить 2? А если вы хотите копать глубже, инструменты профилирования могут предоставить статистику о воспроизведении команд (возникающую из-за расхождения деформации), метриках полосы пропускания / вычисления и т. Д.

Я должен возразить с этим общим заявлением:

Конечно, при использовании графических процессоров любой ценой следует избегать отклонений деформации с помощью операторов if и switch.

Это просто неправда. Способность машины обрабатывать расходящиеся потоки управления на самом деле является функцией, которая позволяет нам программировать ее на более дружественных языках, таких как C/C++, и фактически отличает ее от некоторых других технологий ускорения, которые не предоставляют программисту такой гибкости.

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

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

Теперь для моего ответа на ваши вопросы:

Являются ли накладные расходы (в планировании) расхождения деформации настолько велики, что версия 1) медленнее, чем версия 2?

Это будет зависеть от конкретного уровня ветвления, который на самом деле происходит. В худшем случае, с полностью независимыми путями для 32 потоков, машина будет полностью сериализована, и вы в действительности будете работать на 1/32 пиковой производительности. Разделение потоков по типу дерева двоичных решений не может привести к этому наихудшему случаю, но, безусловно, может приблизиться к нему к концу дерева. Может быть возможным наблюдать более чем 50% -ное замедление этого кода, возможно, 80% или более замедление из-за полного расхождения потоков в конце. Но это будет статистически зависеть от того, как часто происходит дивергенция (то есть зависит от данных). В худшем случае я ожидаю, что версия 2 будет быстрее.

Для версии 2 требуется намного больше ALU, чем для версии 1, и большинство из них тратится на "умножение на 0" (только некоторые из условных выражений оцениваются как 1, а не как 0). Связывает ли это ценные ALU в бесполезных операциях, задерживая инструкции в других деформациях?

float против int может на самом деле помочь здесь, и может быть что-то, что вы могли бы рассмотреть, возможно, изучение. Но во втором случае (для меня) все те же сравнения, что и в первом случае, но с несколькими дополнительными умножениями. В случае с плавающей запятой машина может делать одно умножение на поток за такт, так что это довольно быстро. В случае с int это происходит медленнее, и вы можете увидеть конкретные результаты выполнения команд в зависимости от архитектуры здесь. Я не был бы чрезмерно обеспокоен этим уровнем арифметики. И снова, это может не иметь никакого значения, если ваше приложение ограничено пропускной способностью памяти.

Другой способ выявить все это - написать ядро, которое сравнивает интересующие вас коды, скомпилировать в ptx (nvcc -ptx ...) и сравните инструкции PTX. Это дает гораздо лучшее представление о том, как будет выглядеть код машинного потока в каждом случае, и если вы просто сделаете что-то вроде подсчета команд, вы можете обнаружить небольшую разницу между этими двумя случаями (что должно предпочесть вариант 2 в этом случае),