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 버전을 바이너리 검색을 수행 (약 20ms의.) : 의 길이는 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

커널이하는 일이 무엇인지 완전히 모르겠다. 그렇지만 검색 기준을 만족하는 색인을 하나만 찾고 있다고 가정 할 때 나는 맞는가? 그렇다면 그러한 쿼리를 구조화하고 최적화하는 방법에 대한 몇 가지 지침을 위해 CUDA와 함께 제공되는 축소 샘플을 살펴보십시오.

빠른 포인터하지만 (무엇 당신이하고있는 것은 기본적으로 쿼리에 가장 가까운 인덱스를 줄이기 위해 노력하고있다) : 읽고 매우 느린 글로벌 메모리에 기록의 당신은 엄청 많이 수행하는

. 대신 공유 메모리를 사용해보십시오.

두 번째로 __syncthreads()는 동일한 블록의 스레드 만 동기화하므로 전역 메모리에 대한 읽기/쓰기가 모든 스레드에서 동기화되지는 않습니다 (전역 메모리 쓰기로 인한 대기 시간은 실제로는 그럴 경우)