CUDAバイナリ検索の実装

Question

私はCPUのバイナリ検索をスピードアップしようとしています。残念ながら、GPUのバージョンはCPUのバージョンよりずっと遅いです。おそらく、この問題はGPUには適していないか、何か間違っているのでしょうか？

CPUバージョン（約0.6ms）： 長さ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);  
} 

GPUバージョン（約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のプロセス全体をシリアル化しています。同じワープのすべてのスレッドがブランチ内の同じパスになるように作業を分割したいとします。 CUDA Best Practices Guideの47ページを参照してください。

Answer 2

あなたのカーネルが何をしているのかよく分かっていないのですが、あなたの検索条件を満たす索引を1つだけ探していると仮定しているのでしょうか？もしそうなら、そのようなクエリを構造化して最適化する方法についてのいくつかの指針については、CUDAに付属の削減サンプルを見てください。しかし

いくつかの簡単なポインタ（何あなたがやっていることは、本質的に、あなたのクエリに最も近い指標を削減しようとしている）：

あなたが読み、信じられないほど遅いグローバルメモリへの書き込みの非常に多くを行っています。代わりに共有メモリを使用してみてください。

__syncthreads（）は同じブロック内のスレッドのみを同期するため、グローバルメモリへの読み書きは必ずしもすべてのスレッドで同期されるわけではありません（グローバルメモリ書き込みのレイテンシは実際にはそうであれば）