Как можно уменьшить количество веток / операторов if в моих ядрах?

class tag_option_1 {};
class tag_option_2 {};
class tag_option_3 {};

template<typename T> __device__
int cal_something(int ab, int cd)
{
return -12345; // error value. default case is an error.
};

template<> __device__
int cal_something<tag_option_1>(int ab, int cd)
{
// return something
}

template<> __device__
int cal_something<tag_option_2>(int ab, int cd)
{
// return something
}

template<> __device__
int cal_something<tag_option_3>(int ab, int cd)
{
// return something
}

////////////////////////////////
// version #1:

__global__
void calc_hamiltonian(int * foo, int * bar)
{
unsigned int tid = /* calce thread index*/;

// do something...

if (/* condition */)
{
cal_something<tag_option_1>(foo[tid], bar[tid]);
}
else if(/* condition */)
{
cal_something<tag_option_2>(foo[tid], bar[tid]);
}
else if(/* condition */)
{
cal_something<tag_option_3>(foo[tid], bar[tid]);
}
// no default case.

// do something...
}

////////////////////////////////
// version #2:

// "magical" way to select a version:
// variant is meant to be something like "boost::variant" but implementented without a single "if" statement.
// This "magical" step is meant to be resolved at compile time.
__devcie__
variant <tag_option_1, tag_option_2, tag_option_3>
version_selector(int ab, int cd)
{
// magic happens here.
}__global__
void calc_hamiltonian(int * foo, int * bar)
{
unsigned int tid = /* calce thread index*/;

// do something...

cal_something <version_selector(foo[tid], bar[tid])> (foo[tid], bar[tid]);

// do something...
}

Я в целом согласен с рекомендацией @njuffa не пытаться искусственно искажать свой естественный стиль кодирования, и что вам следует стремиться к производительности (и удобочитаемости и удобству обслуживания), а не к подсчету ветвей в вашем исходном коде. Тем более, что компилятор может заставить их иногда уходить.

Было сказано, что…

Несколько общепринятых «умных» способов уменьшить количество веток (в CUDA и в целом):

1. Подготовьте условия для проверки заранее, желательно во время компиляции.

Объяснение на примере. Версия 1:

 void foo(int* a, bool cond) {
...
for(int i = 0; i < lots; i++) {
if (cond) do_stuff()
else do_other_stuff();
}
...
}

bool cond = check_stuff();
foo(data, cond);

Версия 2:

 void foo(int* a, bool cond) {
...
if (cond) {
for(int i = 0; i < lots; i++) { do_stuff(); }
}
else {
for(int i = 0; i < lots; i++) { do_other_stuff(); }
}
...
}

bool cond = check_stuff();
foo(data, cond);

Версия 3:

 template <bool Cond>
void foo(int* a) {
...
if (cond) {
for(int i = 0; i < lots; i++) { do_stuff(); }
}
else {
for(int i = 0; i < lots; i++) { do_other_stuff(); }
}
...
}

bool cond = check_stuff();
if (cond) foo<true>(data) else foo<false>(data);

И хорошая вещь о версии 3 в том, что пока она выглядит подобно тому, как у него есть ветвь, на самом деле это не так — компилятор принимает только оператор «then» или только оператор «else», но не оба в одной и той же функции.

Переход с версии 1 на версию 2 — это то, что компилятор может быть достаточно любезным для вас; но иногда это не так просто, как в примере, и вы должны позаботиться об этом сами. Переход от версии 2 к версии 3 — это то, что компилятор никогда не сделает для вас.

1.1 Рассмотрим шаблоны ядра на размер блока

Это не всегда — фактически, не часто — полезно, но есть известный пример, приведенный Марком Харрисом в его презентация на оптимизации параллельных сокращений с CUDA. Посмотрите на оптимизацию № 6 на слайдах 24-27. Но не пытайтесь что-то подобное — что уродливо и несколько хрупко — если вы не тщательно приуроченная ваше исполнение, чтобы убедиться, что оно того стоит.

2. Заставьте поведение расходиться в данные а не в поток управления

Версия 1:

 void foo(int* a, int *b) {
...
if (check(a[global_thread_index]) { b[global_thread_index]++; }
}

Версия 2:

 void foo(int* a, int *b) {
...
b[global_thread_index] += check(a[global_thread_index]);
}

(при условии, что проверка возвращает логическое значение или целое число 0 при ошибке и 1 при успехе.)

Здесь я не уверен, что будет делать компилятор CUDA; Кроме того, вы платите штраф за читабельность, написав этот код и, возможно, нарушая «принцип наименьшего удивления». Но вы можете найти менее надуманные примеры.

2.1 Ограничить ветвления значениями выбора тринарной операции

Есть также версия 3:

 void foo(int* a, int *b) {
...
b[global_thread_index] = check(a[global_thread_index]) ? 1 : 0;
}

Теперь у этого все еще есть ветвь — оператор trinary — это просто сокращение для «если», но если вы можете привести свой код к этому состоянию, расхождение будет ограничено одним оператором и каждой ветвью, и, возможно, даже меньше, если компилятору CUDA удается использовать slct PTX оператор:

SLCT

Выберите один исходный операнд на основе знака третьего операнда.

это «оборачивает» семантику ветви в комбинаторную логику для одной инструкции.

Конечно, slct может использоваться компилятором в других случаях; это не ваше дело.

3. Эффективно избегайте ветвей, стремясь к общему консенсусу

(см. также комментарий @ RobertCrovella к тому же эффекту.)

Опять объяснение примером.

 void foo(int* a, int *b) {
...
if (threadIdx.x % 2 == 0) { do_stuff(); }
else { do_other_stuff(); }
...
}

Версия 2:

 void foo(int* a, int *b) {
...
if (threadIdx.x >= blockDim.x / 2) { do_stuff(); }
else { do_other_stuff(); }
...
}

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

Для примера из реальной жизни посмотрите слайды 7-13 в Марк Харрис презентация Я упоминал выше.