OpenCL/CUDA: Algorithme de réduction en deux étapes

Question

La réduction de grands tableaux peut être effectuée en appelant __reduce(); plusieurs fois.

Le code suivant utilise cependant que deux étapes et est documenté here:

Cependant, je suis incapable de comprendre l'algorithme de cette réduction en deux étapes. certains peuvent-ils donner une explication plus simple?

__kernel 
void reduce(__global float* buffer, 
     __local float* scratch, 
     __const int length, 
     __global float* result) { 

    int global_index = get_global_id(0); 
    float accumulator = INFINITY; 
    // Loop sequentially over chunks of input vector 
    while (global_index < length) { 
     float element = buffer[global_index]; 
     accumulator = (accumulator < element) ? accumulator : element; 
     global_index += get_global_size(0); 
    } 

    // Perform parallel reduction 
    int local_index = get_local_id(0); 
    scratch[local_index] = accumulator; 
    barrier(CLK_LOCAL_MEM_FENCE); 
    for(int offset = get_local_size(0)/2; offset > 0; offset = offset/2) { 
     if (local_index < offset) { 
      float other = scratch[local_index + offset]; 
      float mine = scratch[local_index]; 
      scratch[local_index] = (mine < other) ? mine : other; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    if (local_index == 0) { 
     result[get_group_id(0)] = scratch[0]; 
    } 
} 

Il peut également être bien mis en œuvre en utilisant CUDA.

Answer 1

Vous créez des threads N. Le premier thread regarde les valeurs aux positions 0, N, 2 * N, ... Le deuxième thread regarde les valeurs 1, N + 1, 2 * N + 1, ... C'est la première boucle. Il réduit les valeurs length en valeurs N. Puis chaque thread enregistre sa plus petite valeur dans la mémoire partagée/locale. Ensuite, vous avez une instruction de synchronisation (barrier(CLK_LOCAL_MEM_FENCE).) Ensuite, vous avez une réduction standard de la mémoire partagée/locale. Lorsque vous avez terminé, le thread avec l'ID local 0 enregistre son résultat dans le tableau de sortie. En résumé, vous avez une réduction de length à N/get_local_size(0). Vous auriez besoin de faire une dernière passe après l'exécution de ce code. Cependant, cela permet de faire la plus grande partie du travail, par exemple, vous pouvez avoir une longueur ~ 10^8, N = 2^16, get_local_size (0) = 256 = 2^8, et ce code réduit 10^8 éléments en 256 éléments .

Quelles parties ne comprenez-vous pas?