OpenCL: основные вопросы о модели исполнения SIMT

Question

Некоторые концепции и конструкции архитектуры «SIMT» мне до сих пор неясны.

Из того, что я видел и читал, расходящиеся пути кода и вообще if () - довольно плохая идея, потому что многие потоки могут выполняться в режиме lockstep. Что это значит? Как насчет чего-то вроде:

kernel void foo(..., int flag)
{
    if (flag)
        DO_STUFF
    else
        DO_SOMETHING_ELSE
}

Параметр «флаг» одинаков для всех рабочих единиц, и одна и та же ветвь берется для всех рабочих единиц. Теперь, собирается ли GPU выполнить весь код, тем не менее, сериализовать все и в основном по-прежнему брать ветвь, которая не взята? Или это немного умнее и будет выполнять только выполненную ветвь, если все потоки согласуются с принятой ветвью? Что всегда будет здесь.

т.е. сериализация ВСЕГДА происходит или только при необходимости? Извините за глупый вопрос. ;)

Answer 1

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

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

Answer 2

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

Но рассмотрим следующий случай:

kernel void foo(..., int *buffer)
{
    if (buffer[get_global_id(0)])
        DO_STUFF
    else
        DO_SOMETHING_ELSE
}

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

Answer 3

не уверен насчет ati, но для nvidia - это умно.Сериализации не будет, если все потоки в warp будут работать одинаково.