In questo tutorial mostriamo come generare istogrammi tramite Cuda, ovviamente con funzioni __global__ esterne al main e con il cronometraggio dei tempi di calcolo.

- Novità: implementazione della funzione atomicAdd



// CREAZIONE ISTOGRAMMI

#include < iostream >
#include < cstdlib >
#include "sm_11_atomic_functions.h"

using namespace std;

#define NPUNTI 1000000
#define NBINS 256
#define THREADS 256
#define NBLOCKS 32

//////////////////////////////////////////////////////////////////////////////////

__global__ void histogram( unsigned char *buf, unsigned int *his ){
int trid = threadIdx.x;
int tid = blockDim.x*blockIdx.x + trid;
int jump = gridDim.x*blockDim.x;

__shared__ unsigned int temp[NBINS];

temp[trid] = 0;
__syncthreads();

while (tid < NPUNTI){ // scorro i dati a salti di n*k threads
// atomicAdd: operazione che effettua la modifica della variabile temp con il suo indice
// appena accede all'indirizzo lo locka, lo cambia aggiungendogli il numero e poi lo libera
atomicAdd( &temp[buf[tid]], 1 );
tid += jump; // sto aggiornando l´istogramma finale del blocco blockIdx.x
}
__syncthreads();

atomicAdd( &his[trid], temp[trid] ); // sommo tutti gli istogrammi parziali, un blocco alla volta
};

/////////////////////////////////////////////////////////////////////////////////////

int main (void) {

  unsigned char *dev_buf;
  unsigned int *dev_his;
  unsigned char *buf = (unsigned char *)malloc( NPUNTI*sizeof(unsigned char) );
for (int i=0; i
  unsigned int *his = (unsigned int *)malloc( NBINS*sizeof(unsigned int) );

  // faccio partire il cronometro
  cudaEvent_t start,stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start,0);

  // alloco memoria sulla scheda grafica
  cudaMalloc( (void **)&dev_buf, NPUNTI*sizeof(unsigned char) );
  cudaMalloc( (void **)&dev_his, NBINS*sizeof(unsigned int) );
  cudaMemset( dev_his, 0, NBINS*sizeof(unsigned int) );

cout << "programming with cuda c\n";

  // copio su device, lancio il kernel, copio su host
  cudaMemcpy( dev_buf, buf, NPUNTI*sizeof(unsigned char), cudaMemcpyHostToDevice );
  histogram<<>>(dev_buf, dev_his);
  cudaMemcpy( his, dev_his, NBINS*sizeof(unsigned int), cudaMemcpyDeviceToHost );

  for( int i=0; i
cout << i << "\t" << his[i] << "\n";
  };

  // fermo il cronometro
  cudaEventRecord(stop,0);
  cudaEventSynchronize(stop);
  float elapsed;
  cudaEventElapsedTime(&elapsed,start,stop);
  cout << "Time:" << elapsed/1000. << endl;

  // dealloco
  cudaEventDestroy(start);
  cudaEventDestroy(stop);
  cudaFree(dev_buf);
  cudaFree(dev_his);
  free(buf);
  free(his);

return 0;
}