HPC - Prodotto scalare CUDA

Moreno Marzolla moreno.marzolla@unibo.it

Ultimo aggiornamento: 2021-05-14

Familiarizzare con l'ambiente di lavoro

Il server isi-raptor03.csr.unibo.it dispone di tre GPU identiche (NVidia GeForce GTX 1070). Di default viene utilizzata la prima; è però possibile selezionare la scheda da usare sia da programma (es., cudaSetDevice(1) usa la seconda; le GPU sono numerate a partire da zero), sia mediante la variabile d'ambiente CUDA_VISIBLE_DEVICES; ad esempio

    CUDA_VISIBLE_DEVICES=0 ./cuda-stencil1d

esegue il programma cuda-stencil1d sulla prima GPU (default), mentre

    CUDA_VISIBLE_DEVICES=1 ./cuda-stencil1d

lo esegue sulla seconda.

Il comando deviceQuery visualizza le caratteristiche delle GPU.

Prodotto scalare

Modificare il file cuda-dot.cu per calcolare e stampare il prodotto scalare tra due array x[] e y[] di uguale lunghezza n sfruttando la GPU, trasformando la funzione dot() in un kernel. Ricordo che il prodotto scalare s di due array x[] e y[] è definito come

\[ s = \sum_{i=0}^{n-1} x[i] \times y[i] \]

Sono necessarie diverse modifiche alla funzione dot() per sfruttare la GPU. Per questo esercizio si richiede di utilizzare un singolo blocco composto da BLKDIM thread, procedendo come segue:

  1. La CPU alloca sulla GPU un array tmp[] di BLKDIM elementi, oltre ad una copia degli array x[] e y[]. Per allocare l'array tmp[] è possibile utilizzare la funzione cudaMalloc().

  2. La CPU attiva un singolo thread block 1D composto da BLKDIM thread

  3. Il thread \(t\) (\(t = 0, \ldots, BLKDIM-1\)) calcola il valore dell'espressione \((x[t] \times y[t] + x[t + BLKDIM] \times y[t + BLKDIM] + x[t + 2 \times BLKDIM] \times y[t + 2 \times BLKDIM] + \ldots\)) e memorizza il risultato in tmp[t].

  4. Una volta che il kernel termina l'esecuzione, la CPU trasferisce l'array tmp[] dalla memoria del device a quella dell'host, e ne somma il contenuto determinando così il prodotto scalare cercato.

Si rende quindi necessario calcolare il prodotto scalare in due fasi: la prima (passo 3) viene svolta dalla GPU, mentre la seconda (passo 4) viene svolta dalla CPU. La Figura 1 mostra l'assegnazione del calcolo dei prodotti scalari ai thread CUDA, nel caso \(BLKDIM=4\) e \(n=19\).

Figura 1
Figura 1

Il programma deve funzionare correttamente per qualunque valore di \(n\), anche se non è un multiplo di BLKDIM.

Compilare con:

    nvcc cuda-dot.cu -o cuda-dot -lm

Eseguire con:

    ./cuda-dot [len]

Esempio:

    ./cuda-dot

File