Оптимизировать производительность кода, когда нечетные / четные потоки делают разные вещи в CUDA

Question

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

Например:

вектор 1 - это V1 (1), V1 (2), V1 (3), V1 (4)
. вектором 2 является V2 (1) V2 (2) V2 (3) V2 (4)
V1 (1) * V2 (2)
V1 (3) * V2 (4)
V1 (2) * V2 (1)
V1 (4) * V2 (3)

Я написал код Cuda для этого (Pds имеет элементы первого вектора в общей памяти, Nds - второй вектор):

// instead of % 2, checking the first bit to decide if a number
// is odd/even is faster 

if ((tx & 0x0001) == 0x0000)
    Nds[tx+1] = Pds[tx] * Nds[tx+1];
else
    Nds[tx-1] = Pds[tx] * Nds[tx-1];
__syncthreads();

Есть ли еще способ ускорить этот код или избежать расхождения?

Answer 1

Вы должны быть в состоянии устранить ветвь следующим образом:

int tx_index = tx ^ 1; // equivalent to: tx_index = (tx & 1) ? tx - 1 : tx + 1
Nds[tx_index] = Pds[tx] * Nds[tx_index];

Answer 2

Это старый пост, возможно, кто-то посчитает мой ответ полезным. Если в вашем коде tx - threadIdx, то у вас есть расхождение ветвления или деформации. Вы должны избегать расхождения в блоках, потому что он сериализует процесс. Это означает, что сначала будут выполняться потоки с четными индексами, а затем - потоки с нечетными индексами. Если tx - threadIdx, попробуйте изменить свой алгоритм так, чтобы ветвление зависело от blockIdx.