Ядро CUDA с одной веткой работает в 1,5 раза быстрее, чем ядро ​​без ветви

Question

У меня странное снижение производительности на ядре фильтра с разветвлением и без него. Ядро с ветвлением работает в 1,5 раза быстрее, чем ядро ​​без ветвления.

По сути, мне нужно отсортировать пучок лучей, а затем применить ядра взаимодействия. Поскольку сопутствующих данных много, я не могу использовать что-то вроде thrust::sort_by_key() много раз.

Идея алгоритма:

1. Запустите цикл для всех возможных типов взаимодействия (а это пять)
2. В каждом цикле основной поток голосует за свой тип взаимодействия
3. После завершения цикла каждый поток деформации знает о других потоках с таким же типом взаимодействия.
4. Потоки выбирают лидера (по типу взаимодействия)
5. Лидер обновляет таблицу смещений взаимодействий, используя atomicAdd
6. Каждый поток записывает свои данные в соответствующее смещение

Я использовал методы, описанные в этом сообщении Nvidia https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/

Мое первое ядро ​​содержит ветвь внутри цикла и работает около 5 мс:

int active;
int leader;
int warp_progress;
for (int i = 0; i != hit_interaction_count; ++i)
{
  if (i == decision)
  {
    active = __ballot(1);
    leader = __ffs(active) - 1;
    warp_progress = __popc(active);
  }
}

Мое второе ядро ​​использует таблицу поиска из двух элементов, не использует ветвления и работает в течение ~8 мс:

int active = 0;
for (int i = 0; i != hit_interaction_count; ++i)
{
  const int masks[2] = { 0, ~0 };
  int mask = masks[i == decision];
  active |= (mask & __ballot(mask));
}
int leader = __ffs(active) - 1;
int warp_progress = __popc(active);

Общая часть:

int warp_offset;
if (lane_id() == leader)
  warp_offset = atomicAdd(&interactions_offsets[decision], warp_progress);
warp_offset = warp_broadcast(warp_offset, leader);
...copy data here...

Как это может быть? Есть ли способ реализовать такое ядро ​​фильтра, чтобы оно работало быстрее, чем ветвящееся?

UPD: полный исходный код можно найти в filter_kernel cuda_equation / radiance_cuda.cu по адресу https://bitbucket.org/radiosity/engine/src

Answer 1

Я думаю, что это деформация мозга программиста процессора. На CPU я ожидаю повышения производительности из-за устранения ошибки неверного прогнозирования ветвей и ветвей.

Но на GPU нет прогнозирования ветвлений и штрафов, поэтому важны только инструкции.

Сначала мне нужно переписать код на простой.

С филиалом:

int active;
for (int i = 0; i != hit_interaction_count; ++i)
    if (i == decision)
        active = __ballot(1);

Без ветки:

int active = 0;
for (int i = 0; i != hit_interaction_count; ++i)
{
  int mask = 0 - (i == decision);
  active |= (mask & __ballot(mask));
}

В первой версии ~3 операции: compare, if а также __ballot(), Во второй версии ~5 операций: compare, make mask, __ballot(), & а также |=, И в общем коде ~15 операций.

Оба цикла работают в течение 5 циклов. Всего 35 операций в первом и 45 операций во втором. Этот расчет может объяснить снижение производительности.