Implementazione dell'algoritmo di PREFIX SUM in ambiente C Nvidia Cuda

Implementazione dell’operazione di PREFIX SUM inclusiva in ambiente NvidiaCuda.

Obiettivo

Viene fornito un programma C-Cuda in cui inserire il kernel per elaborare il vettore PREFIX SUM inclusiva di un vettore fornito in input. Il kernel sviluppato dovrà utilizzare operazioni di Shuffle. Il programma sarà quindi compatibile con dispositivi Cuda con Compute Capabilities non inferiore a 3.0.

Il problema della PREFIX SUM è stato per tempo considerato come un esempio di come un problema apparentemente seriale possa essere parallelizzato. I due algoritmi più famosi che risolvono tale problema permettono di risolverlo con una complessità di O(n*logn) e O(n).

L’algoritmo sviluppato è quello con complessità maggiore O(logn * n).

Ambiente di sviluppo

L’applicazione è stata sviluppata e testata su una Nvidia Quadro P6000 24GB con Compute Capability di 6.1. L’IDE utilizzato è Atom, il makefile utilizzato è disponibile insieme al codice sorgente consegnato.

Kernel Cuda

Il kernel viene invocato attualmente con un numero di thread per blocco pari a 1024, questo oltre a costituire un limite tecnologico dell’architettura CUDA permette di avere al massimo 32 warp per ogni blocco e quindi poter calcolare in modo efficiente con un solo blocco il valore restituito come risultato dagli altri warp del blocco.

L’algoritmo è composto da 3 parti, in base all’organizzazione gerarchica dei Thread all’interno dell’architettura CUDA.

1. Risoluzione a livello di WARP

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

// sum in warp, <<= is a shift, it doubles the number 
   for (d=1; d<WARPSIZE; d<<=1) {
      others_int = __shfl_up(personal_int,d); // no need for syncthread,
      // shuffle instructions are executed at the same time for all warp
      if (threadId_warp >= d)
      personal_int += others_int;
    }
if (threadId_warp == 31) //last thread of warp store its value for other warps
    vector_block[warpId_block] = personal_int;
__syncthreads(); //data for single warp is updated

Viene risolto il problema della Prefix Sum all’interno di ogni warp utilizzando le istruzioni shuffle. In questo modo i thread comunicano tra loro “scambiandosi” il valore contenuto dai registri, utilizzando un numero di passi logaritmico secondo l’algoritmo naive illustrato sulla documentazione Cuda.

2. Risoluzione a livello di blocco

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

if (warpId_block == 0){
      int other_warp, sum_warp = vector_block[threadId_warp];
for (d=1; d<WARPSIZE; d<<=1) {
        other_warp = __shfl_up(sum_warp,d);
        if (threadId_warp >= d) //maybe not needed
        sum_warp += other_warp;
      }
vector_block[threadId_warp] = sum_warp;
    }
__syncthreads();
    if(warpId_block>0)
      personal_int += vector_block[warpId_block-1];
    data[threadId_absolute] = personal_int;
__syncthreads();

Viene risolto il problema della Prefix Sum all’interno di ogni blocco, partendo dalle soluzioni di ogni warp. All’inizio di questa fase il primo warp (tutti e 32 i thread che ne fanno parte) raccoglie l’ultimo valore di ogni warp del blocco e genera un vettore di Prefix Sum (un vettore di 32 elementi, costituito partendo dai valori dell’ultimo thread di ogni warp, seguendo l’algoritmo illustrato nel punto 1). In questo modo ogni thread può sommare al suo valore calcolato quello del precedente warp, in modo da risolvere la prefix sum all’interno di ogni blocco. La prefix sum effettuata dal primo warp all’inizio di questa fase è stata effettuata utilizzando le operazioni di shuffle, esattamente come per il punto 1. A differenza del punto 1 però, i valori sono stati scambiati dai warp utilizzando un vettore in Shared Memory.

3. Risoluzione globale

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16

if(blockId == 0){ // first block is ready
      if(threadId_block == 0)
      atomicAdd(ready_counter, 1);
      return;
    }
if(blockId != 0 && threadId_block==0){
      int now = atomicAdd(ready_counter, 0);
      while(now < blockId){now = atomicAdd(ready_counter, 0);} // waiting for prev block
    }
    __syncthreads();
personal_int += data[(blockId)*blockDim.x - 1];
    data[threadId_absolute] = personal_int;
__syncthreads();
    if(threadId_block == 0){
      atomicAdd(ready_counter, 1); // this block is ready, next one can start
    }

Non manca che fondere i risultati per ogni blocco. Ogni thread infatti deve solo aggiungere la somma di tutti gli ultimi elementi dei precedenti blocchi al valore calcolato prima per ottenere il risultato finale. Il problema è che la comunicazione fra blocchi non è possibile tramite la veloce Shared Memory, in più non è possibile usare primitive di sincronizzazioni come “__syncthreads()”. Il problema principale è che l’architettura Cuda non specifica l’ordine di esecuzione dei vari blocchi, quindi può avvenire che un blocco i abbia raggiunto questa fase prima del blocco i-1, in questo modo il valore ricercato dai thread del blocco i non è ancora disponibile.

Per sincronizzare i blocchi è stata usata la primitiva “atomicAdd” su un valore di memoria globale. Il primo blocco, non appena conclude la fase 2 incrementa di una unità un contatore globale inizialmente posto a 0 tramite la suddetta primitiva. Il secondo blocco, inizia la fase 3 ponendosi in attesa (il primo thread chiama la atomicAdd sul contatore globale, incrementandolo di 0, gli altri thread del blocco attendono il primo con la “__syncthreads()”) che il contatore diventi uguale ad 1. Una volta che il primo blocco incrementa il contatore, i thread del secondo blocco vengono “sbloccati” e aggiornano il proprio valore recuperandolo dall’ultimo elemento di cui è responsabile il primo blocco. Una volta aggiornati tutti gli elementi del blocco incrementa il contatore sbloccando il terzo blocco di thread. Il meccanismo si ripete fino a quando tutti i blocchi sono stati risvegliati ed hanno svolto la loro somma per ogni thread, restituendo così un vettore corretto che risolve il problema della Prefix Sum.

Repository

La repository con tutto il codice sorgente è disposibile su: https://gitlab.com/Ablablab/multi-many-core-1617