HPC - Odd-even transposition sort

Moreno Marzolla moreno.marzolla@unibo.it

Ultimo aggiornamento: 2021-11-18

A lezione è stato discusso l'algoritmo di ordinamento Odd-Even Transposition Sort. L'algoritmo è una variante di BubbleSort, ed è in grado di ordinare un array di \(n\) elementi in tempo \(O(n^2)\). Pur non essendo efficiente, l'algoritmo si presta bene ad essere parallelizzato. Abbiamo già visto versioni parallele basate su OpenMP e MPI; in questo esercizio viene richiesta la realizzazione di una versione CUDA.

Dato un array v[] di \(n\) elementi, l'algoritmo esegue \(n\) fasi numerate da 0 a \(n–1\); nelle fasi pari si confrontano gli elementi di v[] di indice pari con i successivi, scambiandoli se non sono nell'ordine corretto. Nelle fasi dispari si esegue la stessa operazione confrontando gli elementi di v[] di indice dispari con i successivi (Figura 1).

Figura 1: Odd-Even Sort
Figura 1: Odd-Even Sort

Il file cuda-odd-even.cu contiene una implementazione dell'algoritmo Odd-Even Transposition Sort. L'implementazione fornita fa solo uso della CPU: scopo di questo esercizio è di sfruttare il parallelismo CUDA.

Il paradigma CUDA suggerisce di adottare un parallelismo a grana fine, facendo gestire ad ogni thread il confronto e lo scambio di una coppia di elementi adiacenti. La soluzione più semplice consiste nel creare \(n\) CUDA thread e lanciare \(n\) volte un kernel che sulla base dell'indice della fase (da passare come parametro al kernel), attiva i thread che agiscono sugli elementi dell'array di indice pari o dispari. In altre parole, la struttura di questo kernel è simile a:

__global__ odd_even_step_bad( int *x, int n, int phase )
{
    const int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if ( (idx < n-1) && ((idx % 2) == (phase % 2)) ) {
        cmp_and_swap(&x[idx], &x[idx+1]);
    }
}

Questa soluzione non è però efficiente, perché solo metà dei thread sono attivi durante ogni fase. Realizzare quindi una seconda versione in cui in ogni fase si lanciano \(\lceil n/2 \rceil\) CUDA thread, facendo in modo che tutti i thread siano sempre attivi durante ogni fase. Nelle fasi pari i thread \(0, 1, 2, 3, \ldots\) gestiranno rispettivamente le coppie di indici \((0, 1)\), \((2, 3)\), \((4, 5)\), \((6, 7)\), \(\ldots\), mentre nelle fasi dispari gestiranno le coppie di indici \((1, 2)\), \((3, 4)\), \((5, 6)\), \((7, 8)\), \(\ldots\).

La Tabella 1 illustra la corrispondenza tra l'indice "lineare" dei thread idx, calcolato come nel frammento di codice sopra, e la coppia di indici dell'array che devono gestire.

Tabella 1: corrispondenza tra indice dei thread e dell'array
Indice thread Fasi pari Fasi dispari
0 \((0,1)\) \((1,2)\)
1 \((2,3)\) \((3,4)\)
2 \((4,5)\) \((5,6)\)
3 \((6,7)\) \((7,8)\)
4 \((8,9)\) \((9,10)\)
... ... ...

Per compilare:

    nvcc cuda-odd-even.cu -o cuda-odd-even

Per eseguire:

    ./cuda-odd-even [len]

Esempio:

    ./cuda-odd-even 1024

File