CUDA реализация бинарного поиска

Question

Я пытаюсь ускорить бинарный поиск процессора.К сожалению, версия GPU всегда намного медленнее, чем версия CPU.Возможно, проблема не подходит для графического процессора или я что-то не так делаю?

Версия процессора (около 0,6 мс): с использованием отсортированного массива длины 2000 и двоичный поиск определенного значения

...
Lookup ( search[j], search_array, array_length, m );
...
int Lookup ( int search, int* arr, int length, int& m )
{      
   int l(0), r(length-1);
   while ( l <= r ) 
   {
      m = (l+r)/2;      
      if ( search < arr[m] )
         r = m-1;
      else if ( search > arr[m] )
         l = m+1;
      else
      {         
         return index[m];
      }         
   }
   if ( arr[m] >= search )
      return m;
   return (m+1);      
}

Версия с графическим процессором (около 20 мс): с использованием отсортированного массива длины2000 и бинарный поиск для конкретного значения

....
p_ary_search<<<16, 64>>>(search[j], array_length, dev_arr, dev_ret_val);
....

__global__ void p_ary_search(int search, int array_length, int *arr, int *ret_val ) 
{
   const int num_threads = blockDim.x * gridDim.x;
   const int thread = blockIdx.x * blockDim.x + threadIdx.x;
   int set_size = array_length;

   ret_val[0] = -1; // return value
   ret_val[1] = 0;  // offset

   while(set_size != 0)
   {
      // Get the offset of the array, initially set to 0
      int offset = ret_val[1];

      // I think this is necessary in case a thread gets ahead, and resets offset before it's read
      // This isn't necessary for the unit tests to pass, but I still like it here
      __syncthreads();

      // Get the next index to check
      int index_to_check = get_index_to_check(thread, num_threads, set_size, offset);

      // If the index is outside the bounds of the array then lets not check it
      if (index_to_check < array_length)
      {
         // If the next index is outside the bounds of the array, then set it to maximum array size
         int next_index_to_check = get_index_to_check(thread + 1, num_threads, set_size, offset);
         if (next_index_to_check >= array_length)
         {
            next_index_to_check = array_length - 1;
         }

         // If we're at the mid section of the array reset the offset to this index
         if (search > arr[index_to_check] && (search < arr[next_index_to_check])) 
         {
            ret_val[1] = index_to_check;
         }
         else if (search == arr[index_to_check]) 
         {
            // Set the return var if we hit it
            ret_val[0] = index_to_check;
         }
      }

      // Since this is a p-ary search divide by our total threads to get the next set size
      set_size = set_size / num_threads;

      // Sync up so no threads jump ahead and get a bad offset
      __syncthreads();
   }
}

Даже если я попробую большие массивы, соотношение времени не лучше.

Answer 1

В вашем коде слишком много расходящихся ветвей, поэтому вы по существу сериализуете весь процесс на GPU.Вы хотите разбить работу так, чтобы все потоки в одной и той же деформации проходили по одному пути в ветви.См. Стр. 47 Руководства по оптимальной практике CUDA .

Answer 2

Должен признать, я не совсем уверен, что делает ваше ядро, но прав ли я, предполагая, что вы ищете только один индекс, который удовлетворяет вашим критериям поиска?Если это так, то взгляните на пример сокращения, который поставляется с CUDA, для некоторых указателей о том, как структурировать и оптимизировать такой запрос.(То, что вы делаете, по сути пытается уменьшить ближайший индекс к вашему запросу)

Некоторые быстрые указатели, хотя:

Вы выполняете очень много операций чтения и записи в глобальную память, чтоневероятно медленныйПопробуйте вместо этого использовать разделяемую память.

Во-вторых, помните, что __syncthreads () синхронизирует только потоки в одном и том же блоке, поэтому ваши операции чтения / записи в глобальную память не обязательно синхронизируются между всеми потоками (хотя задержка от вашего глобальногозапись в память может фактически заставить это казаться, как будто они делают)