Cuda Control расхождение

Question

говорят, что у меня есть 3 массива общей памяти: a b c. Я не уверен, приведет ли следующее расположение нитей к расхождению в контроле

if (threadIdx < 64)
{
    if (threadIdx == 1)
        for (int i = 0; i < N; i++)
            c += a[threadIdx]*a[threadIdx];
    else
        for (int i = 0; i < N; i++)
            c += a[threadIdx]*b[threadIdx];
}

если да, то насколько плохо это повлияет на производительность? Есть ли эффективный способ справиться с проблемой? спасибо

Answer 1

В зависимости от размеров вашего блока первое условие threadIdx.x < 64 (обратите внимание на .x) может вообще не вызывать расхождение. Например, если у вас есть блок с измерениями (128,1,1), то первые два перекоса (группы из 32 потоков, которые выполняются в шаге блокировки) войдут в блок if, а последние два обойдут его. Поскольку вся деформация идет в одну или другую сторону, дивергенции нет.

Условное выражение, подобное threadIdx.x == 1, вызовет расхождение, но оно будет иметь очень скромные затраты. Действительно, во многих случаях CUDA сможет реализовать условное выражение с помощью одной инструкции. Например, такие операции, как min, max и abs, как правило, будут реализованы с помощью одной инструкции и не вызовут абсолютно никаких расхождений. О таких инструкциях можно прочитать в Руководстве по PTX .

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

if (threadIdx.x % 4 == 0)
  // do expensive operation
else if (threadIdx.x % 4 == 1)
  // do expensive operation
else if (threadIdx.x % 4 == 2)
  // do expensive operation
else
  // do expensive operation

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

Имейте в виду, что расхождение потоков является гораздо меньшей проблемой, чем (1) выбор алгоритма высокого уровня и (2) локальность / объединение памяти. Очень немногие программисты CUDA должны когда-либо беспокоиться о расхождениях в ваших примерах.

Answer 2

Если имеется более одного потока на блок, я бы ожидал расхождения в одной деформации каждого блока (какой бы блок не содержал поток 1).

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

if (threadIdx.x < 64)
{
    //this conditional might diverge
    if (threadIdx.x == 1)
        ptr = a;
    else
        ptr = b;

    //but obviously this part will not
    for (int i = 0; i < N; i++)
        c += a[threadIdx]*ptr[threadIdx];
}