Снижение или атомный оператор на неизвестных глобальных индексов массива

Question

У меня есть следующий алгоритм:

__global__ void Update(int N, double* x, double* y, int* z, double* out) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    if (i < N) 
    { 
     x[i] += y[i]; 
     if (y[i] >= 0.) 
     out[z[i]] += x[i]; 
     else 
     out[z[i]] -= x[i]; 
    } 
} 

Важно отметить, что из меньше х. Скажем, x, y и z всегда одного и того же размера, скажем 1000, а out всегда меньше, скажем, 100. z - это индексы, которые соответствуют каждому из x и y.

Это все находка, кроме обновлений. Там могут быть конфликты между потоками, поскольку z не содержит только уникальные значения и имеет дубликаты. Поэтому в настоящее время это реализовано с атомными версиями atomicAdd и вычитанием с помощью сравнения и свопинга. Это, очевидно, дорого и означает, что мое ядро ​​занимает 5-10 раз дольше.

Я хотел бы уменьшить это, но единственный способ, которым я могу это сделать, - это для каждого потока иметь свою собственную версию (которая может быть большой, 10000+, X 10000+ потоков). Это означало бы, что я установил 10000 double [10000] (возможно, в shared?), Чтобы вызвать мое ядро, а затем суммировать эти массивы, возможно, в другом ядре. Неужели должен быть более элегантный способ сделать это?

Возможно, стоит отметить, что x, y, z и out находятся в глобальной памяти. Поскольку мое ядро ​​(у меня есть другие подобные) очень просто, я не решил копировать бит в общий (nvvp на ядре показывает равные вычисления и память, поэтому я думаю, что не так много производительности при добавлении накладных расходов на перенос данных из глобальное разделение и возвращение назад, любые мысли?).

Answer 1

Метод 1:

1. Построить набор "операций". Поскольку у вас есть только одно обновление для потока, вы можете легко создать запись транзакции фиксированного размера, по одной записи на поток. Предположим, у меня есть 8 потоков (для простоты представления) и какое-то произвольное количество записей в моей таблице out. Давайте предположим, что мои 8 потоков хотел сделать 8 сделок, как это:

   thread ID (i): 0  1  2  3  5  6  7 
   z[i]:   2  3  4  4  3  2  3 
   x[i]:   1.5 0.5 1.0 0.5 0.1 -0.2 -0.1 
   "transaction": 2,1.5 3,0.5 4,1.0 4,0.5 3,0.1 2,-0.2 3,-0.1 
   

2. Теперь сделать sort_by_key по сделкам, чтобы расположить их в порядке z[i]:

   sorted:   2,1.5 2,-0.2 3,0.5 3,-0.1 3,0.1 4,1.0 4,0.5 
   

3. Теперь сделать операцию reduce_by_key по сделкам:

   keys:   2  3  4  
   values:   1.3 0.5 1.5 
   

4. Теперь обновление out[i] по клавишам:

     out[2] += 1.3 
         out[3] += 0.5 
         out[4] += 1.5 
   

thrust и/или cub может быть предварительно построенные варианты сортировки и уменьшения операций.

Метод 2:

Как вы говорите, у вас есть массивы x, y, z и out в глобальной памяти.Если вы собираетесь использовать z, который является «отображение» несколько раз, вы можете захотеть изменить (группа) или сортировки массивов в порядке z:

index (i): 0  1  2  3  4  5  6  7 
     z[i]: 2  8  4  8  3  1  4  4 
     x[i]: 0.2 0.4 0.3  0.1 -0.1 -0.4  0.0 1.0 

группы по г [я]:

index (i): 0  1  2  3  4  5  6  7 
     z[i]: 1  2  3  4  4  4  8  8 
     x[i]:-0.4 0.2 -0.1  0.3 0.0  1.0  0.4 0.1 

Это или его вариант позволит вам избежать повторной операции сортировки в методе 1 (опять же, если вы повторно используете один и тот же вектор отображения).