CUDA二进制搜索执行

我想加快CPU二进制搜索。不幸的是，GPU版本总是比CPU版本慢得多。也许问题不适合GPU或者我做错了什么？CUDA二进制搜索执行

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(); 
    } 
}

即使我尝试更大的阵列，时间比例并没有更好的。

来源

2011-08-16 Izidor

简单的二进制搜索并不完全适合GPU操作。这是一个无法并行化的串行操作。但是，您可以将数组拆分为小块，然后在每个块上执行二进制搜索。创建X块，确定哪些可能包含X并行线程中的变量。抛出所有，但候选人，进一步细分，等等... –

您可能想要检查推测二进制搜索在http://wiki.thrust.googlecode.com/hg/html/group__binary__search.html – jmsu

你的代码中有太多不同的分支，所以你基本上是序列化GPU上的整个过程。您想分解工作，以便同一个warp中的所有线程在分支中采用相同的路径。请参阅CUDA Best Practices Guide的第47页。

来源

2011-08-16 16:04:22 SpaceghostAli

我使用了数组2000元素。并使用编号为395的二进制搜索的CPU版本。在我的PC上它仅花费0.000933ms。为了测试，我创建了内核<<<2000,1> >>，并且内核完全没有做任何事情：__global__ void Search（） {td = threadIdx.x + blockIdx.x * blockDim.x; if（tid <2000） { } }并且仅仅调用它就花了0.034704毫秒。从这个结果我真的想知道是否有意义使用CUDA来使事情变得更快。或者我做错了什么...... – Izidor

这真的就是这样，CUDA作为一些需要一些时间的开销，但是当你正在做的事情，例如，在CPU上10秒，并且GPU可以使它即使有0.03秒的开销，你会更喜欢哪一种？ CUDA绝对有效，但如果在CPU上它已经非常快，它可能不值得。 – jmsu

谢谢。我将尝试在GPU上传输更多PC上的工作，我希望这能够弥补现有的开销。我认为没有CPU和GPU之间的内存拷贝，开销已经很小，但显然不是。我也将检查“推力”二进制搜索。 – Izidor

我必须承认我不完全确定你的内核是做什么的，但是我认为你正在寻找一个满足你的搜索条件的索引吗？如果是这样，那么查看一下CUDA附带的简化示例，了解如何构建和优化这样的查询。（什么你正在做的基本上是试图最接近的索引减少到您的查询），尽管

一些简单的指针：

的读取和写入全局存储器，这是令人难以置信的慢你正在执行一个可怕的很多。尝试使用共享内存。其次请记住，__syncthreads（）只能同步同一个块中的线程，因此您对全局内存的读取/写入不一定会在所有线程中同步（尽管全局内存写入的延迟可能实际上使其显示为如果他们这样做）

来源

2012-09-09 19:15:07

CUDA二进制搜索执行

回答

相关问题