Используйте CUDA для эффективного вычисления позиций отсортированного массива, где изменяется элемент

Допустим, у нас есть этот отсортированный массив

     0 1 1 1 1 2 2 2 2 2 3 10 10 10

Я хотел бы эффективно найти позиции, где элемент меняется. Например, в нашем массиве позиции следующие:

    0 1 5 10 11

Я знаю, что есть несколько библиотек (Thrust), которые могут достичь этого, однако я хотел бы создать свою собственную эффективную реализацию для образовательных целей.

Вы можете найти весь код здесь: http://pastebin.com/Wu34F4M2

Это также включает в себя проверку.

Ядро имеет следующую функцию:

__global__ void findPositions(int *device_data,
int totalAmountOfValuesPerThread, int* pos_ptr, int N){

int res1 = 9999999;
int res2 = 9999999;
int index = totalAmountOfValuesPerThread*(threadIdx.x +
blockIdx.x*blockDim.x);
int start = index; //from this index each thread will begin searching
if(start < N){ //if the index is out of bounds do nothing
if(start!=0){ //if start is not in the beginning, check the previous value
if(device_data[start-1] != device_data[start]){
res1 = start;
}
}
else res1 = start; //since it's the
//beginning we update the first output buffer for the thread
pos_ptr[index] = res1;

start++; //move to the next place and see if the
//second output buffer needs updating or not

if(start < N && device_data[start] != device_data[start-1]){
res2 = start;
}

if((index+1) < N)
pos_ptr[index+ 1] = res2;
}
}

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

1. device_data хранит все числа в массиве
2. totalAmountOfValuesPerThread в этом случае — это общее количество значений, с которыми каждый поток должен будет работать
3. pos_ptr имеет ту же длину, что и device_dataкаждый поток записывает результаты буфера в этот device_vector
4. N общее количество чисел в device_data массив

В выходных буферах называется res1 а также res2 каждая нить либо сохраняет позицию, которая не была найдена ранее, либо оставляет ее как есть.

Пример:

  0   <---- thread 1
1
1   <---- thread 2
1
2   <---- thread 3
2
3   <---- thread 4

Выходные буферы каждого потока, предполагая, что большое число 9999999 inf было бы:

  thread1 => {res1=0, res2=1}
thread2 => {res1=inf, res2=inf}
thread3 => {res1=4, res2=inf}
thread4 => {res1=6, res2=inf}

Каждый поток будет обновлять pos_ptr device_vector поэтому этот вектор будет иметь в результате следующее:

  pos_ptr =>{0, 1, inf, inf, 4, inf, 6, inf}

После доработки ядра я сортирую вектор с помощью библиотеки Thrust и сохранить результаты внутри вектора хоста под названием host_pos, Итак host_pos вектор будет иметь следующее:

  host_pos => {0, 1, 4, 6, inf, inf, inf, inf}

Эта реализация ужасна, потому что

1. В ядре создано много веток, поэтому произойдет неэффективная обработка переноса
2. Я предполагаю, что каждый поток работает только с 2 значениями, что очень неэффективно, потому что создано слишком много потоков
3. Я создаю и передаю device_vector который так же велик, как вход, а также находится в глобальной памяти. Каждый поток обращается к этому вектору для записи результатов, это очень неэффективно.

Вот тестовый пример для ввода размера 1 000 000 когда имея 512 темы в каждом блоке.

     CPU time: 0.000875688 seconds
GPU time: 1.35816 seconds

Еще один тест для ввода размера 10 000 000

     CPU time: 0.0979209
GPU time: 1.41298 seconds

Обратите внимание, что версия процессора стала почти в 100 раз медленнее, в то время как GPU показал почти то же время.

К сожалению, у моего GPU недостаточно памяти, поэтому давайте попробуем 50 000 000

     CPU time: 0.459832 seconds
GPU time: 1.59248 seconds

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

Какой дизайн вы бы предложили, чтобы мой алгоритм работал быстрее? К сожалению, я не могу придумать ничего лучшего.

заранее спасибо