In questo tutorial vediamo come implementare la somma di vettori tramite l'algoritmo in CUDA. Notiamo il cronometro per misurare la velocità del programma.


#include < iostream >
using namespace std;
#include < cstdlib >

#define N 10000

// ogni kernel viene dichiarato con __global__
// gira sulla gpu
__global__ void somma_vec( int *a, int *b, int *c){ // i singoli elementi del vettore sono int!

  int tid = blockIdx.x; // Idx è l'indice del blocco -> preso solo lungo x, struttura lineare

// sono nel blocco k? ok, sommo le k-esime di a e b, e le metto nella k-esima componente di c!
// N è il numero di componenti del vettore
// questo controllo mi aiuta a gestire i threads in eccesso
  if(tid < N)
    c[tid]=a[tid]+b[tid]; 
}
// il loop è sostituito dal fatto che ogni blocco fa contemporaneamente la somma delle
// componenti corrispondenti al numero del blocco
// 65.000 circa = numero max di blocchi della scheda: se il vettore ha più componenti avrà segmentation fault

int main( void ){

  int a[N],b[N],c[N];
  int *dev_a, *dev_b, *dev_c;

// cronometro (gestione del tempo di calcolo in ms)
  cudaEvent_t start1,stop1;
  float time1;

  for( int i=0; i
    a[i] = 2*i;
    b[i] = -i;
  };

  // la possibilità di misurare la durata di un evento viene resa possibile
  // dall'utilizzo di tag, che possono essere associati a due istanti temporali
  // e al relativo utilizzo della GPU
  // se non creo gli eventi le menorie non vengono allocate e non posso fare il conto temporale
  cudaEventCreate(&start1);
  cudaEventCreate(&stop1);

// la memoria richiesta sul device deve essere allocata:
// dev_c puntatore a intero, sulla scheda viene allocata memoria pari al
// secondo argomento di cudaMalloc associata all'indirizzo di dev_c
// (void **) pointer a pointer
  cudaMalloc( (void **)&dev_a, N*sizeof(int)  ); // vettori a N componenti
  cudaMalloc( (void **)&dev_b, N*sizeof(int)  ); // vettore a, vettore b
  cudaMalloc( (void **)&dev_c, N*sizeof(int)  ); // risultato vettore c

  //da questo istante misuriamo il tempo trascorso sulla GPU
  cudaEventRecord(start1,0);

  // i dati contenuti in a e b vengono copiati dall'HOST al DEVICE
  // e vengono associati a dev_a e a dev_b
  cudaMemcpy( dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice); // copio da cpu a gpu
  cudaMemcpy( dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice); // l'info copiata è lunga N*sizeof(int)

// il kernel viene invocato  chiedendo che ci siano N blocchi ciascuno 
// contenente un singolo thread
// il risultato viene associato a dev_c
  somma_vec<<>>(dev_a, dev_b, dev_c); // N = 10.000 < 65.000 (N blocchi, 1 thread per blocco)

// il risultato è puntato su dev_c

// nella variabile c viene copiato dal DEVICE all'HOST
// quanto si trova sul device a partire da dev_c
// per una lunghezza pari a N*sizeof(int)

  cudaMemcpy( c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost); // copio dalla gpu alla cpu

// ora posso permettermi di visualizzare su file o su terminale i risultati del calcolo
//   for( int i=0; i
//     cout << "c["<< i <<"]=" << a[i] << b[i] <
//   };

  // libero la memoria di cuda
  cudaFree( dev_a);
  cudaFree( dev_b);
  cudaFree( dev_c);

  // il contatore finale viene misurato
  cudaEventRecord(stop1,0);

  // si chiede che la GPU abbia completato le sue operazioni prima di
  // restituire il valore del contatore stop1
  cudaEventSynchronize(stop1);

  // la differenza tra start1 e stop1 corrisponde al tempo trascorso
  // misurato in millisecondi
  cudaEventElapsedTime(&time1, start1, stop1);
  cout << "time1=" << time1 <<"\n";

  cudaEventDestroy(start1);
  cudaEventDestroy(stop1);

  return 0;
}