поразрядное исключающее ИЛИ

Имеется пример редукции(сложения) массива с участком кода:

  1. for ( i = 0; n >= BLOCK_SIZE; n /= (2*BLOCK_SIZE), i++ ){
  2.         dim3 dimBlock (BLOCK_SIZE, 1, 1);
  3.         dim3 dimGrid (n / (2*dimBlock.x), 1, 1);
  4.         reduce4<<<dimGrid,dimBlock>>>(adev[i], adev[i^1]);
  5. }

объясните зачем тут поразрядное исключающее или i^1.
я понимаю что:
i | i^1
0|1
1|0
2|3
3|2
4|5
5|4

Forums: 

Обычно редукция выполняется

Обычно редукция выполняется примерно так:
тредом грузится порция данных из глобальной памяти в разделяемую.
Далее блок потоков делает редукцию в разделяемой памяти
Результат для блока - одно число, записывается в массив в глобальной памяти.
Дальше нужно за-reduce'ить получившийся массив и так далее.

Итак что нам нужно? входные данные, буфер1, буфер2(если исходные данные нельзя изменять)

на первом шаге reduce'им входые данные в буфер1
потом данные из буфера1 в буфер2
потом из буфера2 в буфер1 и т.д.

Я думаю вы привели код отсюда: http://systemsinside.kiev.ua/node/196?page=show
смотрим:
int * adev [2] = { NULL, NULL };

Мне кажется автор этого кода хотел сделать что-то типа (adev[i & 1], adev[(i^1)&1]);
то есть
0 1
1 0
0 1
1 0
и т.д. - то есть меняем буферы местами

Скорей всего он тестировал это на небольших объёмах данных, где цикл у него выполнялся два раза.
При размере блока 512, чтобы цикл выполнился больше двух раз, n должно быть как минимум больше миллиона.