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 n'utilise cependant que deux étapes et est documenté aquí :

Cependant, je ne parviens pas à comprendre l'algorithme de cette réduction en deux étapes. Quelqu'un peut-il 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 N fils. Le premier fil examine les valeurs aux positions 0, N, 2*N, ... Le deuxième thread examine les valeurs 1, N+1, 2*N+1, ... C'est la première boucle. Elle réduit length en N valeurs.

Ensuite, chaque thread sauvegarde sa plus petite valeur dans la mémoire partagée/locale. Ensuite, vous avez une instruction de synchronisation ( barrier(CLK_LOCAL_MEM_FENCE) .) Vous avez alors 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 somme, vous avez une réduction de length a N/get_local_size(0) valeurs. Vous devrez effectuer un dernier passage après l'exécution de ce code. Cependant, cela permet de faire la plupart du travail, par exemple, vous pourriez avoir 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 ?