Implémentation de la recherche binaire CUDA

Question

J'essaie d'accélérer la recherche binaire du CPU. Malheureusement, la version GPU est toujours beaucoup plus lente que la version CPU. Peut-être que le problème ne convient pas au GPU ou que je fais quelque chose de mal?

version CPU (environ de 0,6 ms.): utilisant le tableau trié de longueur 2000 et faire la recherche binaire pour valeur spécifique

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

version GPU (environ 20ms.): en utilisant le tableau trié de longueur 2000 et faire une recherche binaire pour valeur spécifique

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

Même si j'essaie de plus grands tableaux, le rapport de temps n'est pas mieux.

Answer 1

Vous avez trop de branches divergentes dans votre code, ce qui signifie que vous êtes en train de sérialiser l'ensemble du processus sur le GPU. Vous voulez décomposer le travail de sorte que tous les fils d'une même chaîne prennent le même chemin dans la branche. Voir page 47 du CUDA Best Practices Guide.

Answer 2

Je dois avouer que je ne suis pas tout à fait sûr de ce que fait votre noyau, mais ai-je raison de supposer que vous recherchez un seul index répondant à vos critères de recherche? Si c'est le cas, jetez un œil à l'exemple de réduction fourni avec CUDA pour obtenir des conseils sur la façon de structurer et d'optimiser une telle requête. (Ce que vous faites est essentiellement tente de réduire l'indice le plus proche de votre requête)

Quelques indications rapides si:

Vous effectuez un très grand nombre de lectures et d'écritures à la mémoire globale, ce qui est incroyablement lent. Essayez d'utiliser la mémoire partagée à la place. Souvenez-vous ensuite que __syncthreads() ne synchronise que les threads dans le même bloc, donc vos lectures/écritures dans la mémoire globale ne seront pas nécessairement synchronisées sur tous les threads (bien que la latence de vos écritures de mémoire globales s'ils le font)