Abbiamo due modi per vivere l'oratorio feriale: come il nonno e il papà, oppure come il ragazzo.

Il nonno e il papà sanno fare benissimo il loro lavoro, e hanno ben chiaro lo scopo, cioè far cambiare fase alla luna. Ma, come si vede nel corto, non se lo ricordano, infatti solo alla fine lo capiremo. Il loro lavoro diventa quindi meccanico, entrambi cercano di educare il ragazzo imponendogli la loro visione senza aprire lo sguardo per davvero (infatti hanno sempre gli occhi coperti, chiusi nella loro routine). Il ragazzo cerca di imparare ma non riesce a capire quello che deve fare: non ha una guida.

L'alternativa, il comportamento del ragazzo. Si rende conto che svolgere il suo compito senza aver chiaro lo scopo non gli basta (lui non sa ancora che con il suo lavoro cambierà di fase alla luna).

Dove avviene il cambiamento? Quando un fatto straordinario entra nella loro vita, rappresentato nel corto come una stella più grande delle altre che si schianta sulla luna. Qui si vede ancora l'alternativa: il nonno e il papà non si lasciano colpire, cercano di spostarla come se fosse come le altre stelline. Il ragazzo invece rimane stupefatto e non si comporta come loro, bensì lascia che quello stupore li colpisca ponendosi come tramite. In quel momento il nonno ed il padre aprono per la prima volta gli occhi, e da lì in poi li avranno sempre aperti. Solo così potranno fare memoria del motivo per cui ogni sera vanno sulla luna.

Vi auguro che quell'evento straordinario che per la prima volta ci ha colpiti e senza il quale non saremmo qui ora possa ogni giorno riaccadere per farci ricordare perchè siamo educatori; magari non accadrà la stessa cosa che è successa a Maria grazie all'Angelo, ma quell'incontro si presenterà in fatti quotidiani come lo sguardo di un bambino, l'abbraccio di un amico o le parole del don. Di sicuro, se come il ragazzo ci lasceremo colpire, tutto sarà più bello e potremo anche noi dire "Magnificat anima mea Dominum". Magari neanche noi capiremo subito la nostra chiamata, come nel corto in cui si capsice solo alla fine che il loro lavoro serve per cambiare le fasi della luna. Magari la fatica non verrà da subito capita, ma alla fine lo scopriremo, e scopriremo di essere di Cristo anche lì.

Noi non ci accontentiamo di vedere la bellezza, anche se il Cielo sa che gran dono sia questo. Noi vogliamo qualcos'altro, che è difficile descrivere a parole: vogliamo sentirci uniti alla bellezza che vediamo, trapassarla, riceverla dentro di noi, immergerci in essa, diventarne parte. Ecco perché abbiamo popolato l'aria, la terra e l'acqua di dei e dee, e ninfe ed elfi che, dato che a noi non è possibile, possano almeno loro, queste proiezioni di noi stessi, godere ritrovando in se stessi quella bellezza, quella grazia, e quel potere di cui la Natura è immagine. Ecco perché i poeti ci narrano tante meravigliose menzogne. Essi parlano come se il Vento dell'Occidente potesse davvero soffiare dentro un'anima umana: ma non è vero. Ci dicono che "la Bellezza nata da un sussurro" passerà in un volto umano: ma non lo farà. Non per ora. E se prendiamo sul serio l'immagine della Scrittura, se crediamo che Dio un giorno ci darà davvero la Stella del Mattino e ci farà indossare lo splendore del sole, allora possiamo credere che gli antichi miti, come la poesia moderna, pur così falsi nel loro significato storico, siano tanto vicini alla verità quanto la profezia. In questo momento noi ci troviamo all'esterno del mondo, dalla parte sbagliata della porta. Possiamo percepire la freschezza e la purezza del mattino, senza però che queste ci possano rendere freschi puri come loro. Non possiamo penetrare lo splendore che vediamo. Ma tutte le foglie del Nuovo Testamento sussurrano frusciando che non sarà sempre così. Un giorno, a Dio piacendo, riusciremo ad entrare. Quando le anime umane saranno diventate così perfette nella loro obbedienza volontaria da uguagliare le creature inanimate nella loro obbedienza senza vita, allora potranno riscoprirsi della medesima gloria della natura, anzi, di quella Gloria ben più grande di cui la natura non è che un primo abbozzo. La Natura è mortale: noi le sopravviveremo. Quando tutti i soli e tutte le nebulose saranno tramontati, ognuno di voi sarà ancora vivo. La Natura non è che un'immagine, un simbolo, ma è il simbolo che la Scrittura mi invita ad usare. Siamo invitati ad entrare attraverso la Natura, oltrepassandola, fino a raggiungere quello splendore che essa è in grado di riflettere solo in parte. E una volta dentro, oltrepassata la Natura, potremo mangiare dell'albero della vita.

C. S. Lewis

Il bene di tutti

Gli affreschi del Buon Governo nel palazzo pubblico di Siena
Incontro con Mariella Carlotti
Concorezzo – 21 febbraio 2013



Abbiamo invitato Mariella, già venuta a parlare qui a Concorezzo dell'annuncio a Maria, incontro che ha lasciato il segno sulla nostra comunità. Il bene comune sarà l’argomento di stasera: come ci poniamo di fronte alla situazione delle elezioni? Mariella ci farà vedere con questo tema fu affrontato da Lorenzetti nel suo famoso ciclo di affreschi nel palazzo pubblico di Siena.

Mariella: dico sempre ai miei alunni una cosa, quando devo iniziare un corso di storia in una classe nuova, cioè di immaginare di essere a bordo di un astronave in vista della terra. Quanti continenti vedete? Quattro o cinque? Qual è il quinto? Non c’è  nella geografia: il quinto è una penisola terrestre dell’Asia che chiamiamo Europa (qui capisco chi avrà il 6…). La domanda successiva deve essere: che cosa è accaduto qui? Non è un continente per  la geografia  ma per la storia sì (e qui capisco chi avrà 8…). L’Europa è un continente perché, nel corso della storia, si è fatta una precisa idea del lavoro, della politica e della società: se non riprendiamo coscienza di queste cose siamo solo una penisola dell’Asia. Stasera parliamo della concezione di politica così come nella nostra civiltà si è espressa, e come l’ha espressa Lorenzetti nei suoi affreschi. Siamo a Siena: secondo la leggenda Siena è stata fondata dai figli di Remo che scapparono da Roma dopo aver rubato la statua dei gemelli, su due cavalli (uno bianco e uno nero). Prese il nome da uno dei due fratelli. I fiorentini dicono invece che fu fondata da Brenno, un generale barbaro, che lì lasciò tutti i suoi soldati vecchi e buoni a nulla, e che divenne città solo perché una governante molto generosa con tutti andò a convincere un cardinale (quindi sarebbero dei vecchi, abbandonati, buoni a nulla e figli di…). Siena invece nasce come un villaggio etrusco, che non doveva diventare città perché non aveva acqua: niente laghi, mari o fiumi. Quando i Longobardi aprirono la via Francigena, che attraversava proprio Siena, iniziò la crescita; dopo il mille divenne una città di una certa grandezza, diventando la padrona di tutta la toscana meridionale, e quindi entrando in contrasto con Firenze. Contrasto che finirà dopo 300 anni con la conquista fiorentina senese nel 1550. La prima battaglia di alto valore storico fu la battaglia di Montaperti (4 settembre 1260): Firenze strinse d’assedio Siena con una milizia che superava il doppio di quella senese. C’erano due alternative per Siena: arrendersi o combattere. Il governo senese si ritirò in una chiesa mentre il popolo all’esterno pregava affinché trovasse una soluzione: decisero di combattere, con pochissime possibilità di vincere, ma prima della battaglia proclamarono la Madonna regina di Siena, consacrando a Lei la città. Diedero alla Madonna le chiavi della città, con rogito notarile. Fino ad oggi mai venne un re a Siena: legalmente ci fu sempre la regina.


E, nella battaglia, le milizie senesi distrussero le milizie fiorentine, senza un senso logico, senza spiegazione razionale dato il numero esiguo di combattenti. La vittoria di Montaperti è ricordata ancora oggi: sulle auto senesi ci sono ancora adesivi che recitano “4 settembre 1260 Montaperti c’ero anch’io”. Nei successivi novanta anni Siena divenne guelfa, e il potere venne dati al governo dei nove: due mesi di incarico, senza potersi ricandidare (proprio come oggi…). La politica fu legata al benessere di tutti i cittadini. Ricordo che a Siena tutto è dedicato a Maria: Piazza del Campo è il suo manto, per Maria si corre il palio, il capodanno è il 25 marzo, giorno dell’Annunciazione. In un famoso quadro la Madonna offre Siena a Gesù. La biccherna a Siena era la ragioneria della città: gli sposati non potevano toccare il denaro, quindi erano i monaci a fare quel lavoro. Dico sempre che quando in una città sono belli sia la cattedrale che il palazzo è una cosa normale, ma quando sono bellissimi anche i libri contabili, allora la bellezza ha preso tutto. C’erano perfino dei vigili urbani della bellezza! Così si arriva nel 1310, 50 anni da Montaperti, il Giubileo. Furono fatte tre grandi opere pubbliche: la traduzione del corpo delle leggi (costituzione) in volgare (la prima costituzioni in volgare della storia d’Italia), la lastricatura della Piazza del Campo (nove spicchi, a simboleggiare il governo dei nove), e infine il palazzo di Maria, la regina di Siena (il palazzo di Piazza del Campo). In quei due mesi in cui i nove erano in carica non potevano uscire da quel palazzo per dedicarsi totalmente al governo e al bene della città (e per non essere corrotti, proprio come oggi…). Erano strapagati, così non aveva senso neanche pensare di corromperli. Avevano la messa alle 6 della mattina, tutti i giorni, governo e parlamento (composto da 200 persone). Nella sala in cui si riunivano c'erano le grandi battaglie della città affrescate sui muri, e tutti i più grandi santi sotto di esse: per il medioevale la vita è una cosa drammatica che però poggia sulla certezza del rapporto con Dio. La Madonna, affrescata nella parere centrale, è ciò che unisce il governo a Dio.



La sala del parlamento


Le parole che dice Gesù (nella pergamena che ha in mano, incollata all’affresco) sono le parole d’inizio del libro della dal Sapienza: “Amate la giustizia voi che governate la terra”. Ai Patroni di Siena la Madonna rivolge discorso in terzina dantesca, 1315: ciò dimostra che la Divina Commedia era conosciuta anche con Dante vivente. Questi scritti li mise Cino da Pistoia, amico di Dante. Discorsi: “chi più parla peggio più è lodato” (profetici…) ecc. “Ma le vostre preghiere le esaudirò tutte” ecc. Vent’anni dopo Lorenzetti affresca una sala che comunica con la sala precedente con il ciclo “Il buon governo e i suoi effetti” e “Il mal governo e i suoi effetti”. Inizialmente erano gli affreschi del “Bene comune”, il nome “Buon governo” è dato nel rinascimento. Se la società fa lo stato, è il bene comune. Se lo stato fa la società, allora è buon governo. Qui già si capiscono tante cose della politica di oggi. I problemi della contemporaneità sono innanzitutto il fatto che nessuno dice che le due sale sono comunicanti, quella con la Madonna, tramite tra Dio e il Governo, e quella con il Bene comune. Il secondo problema della contemporaneità è che Lorenzetti spiegò gli affreschi con 62 versi, conscio che l’interpretazione dei posteri sarebbe stata fuorviante: questi versi non sono mai riportati in nessun libro di storia né di arte! Sono iscrizioni dimenticate dagli storici, perché dicono cose troppo precise e scandalose per l’uomo moderno.



Il bene comune

“Questa santa virtù là dove regge indice ad unità li animi molti” è modo per dire che gli affreschi sono legati a quelli di prima. Ma la giustizia non è quella di Ingroia, è quella di Dio. Giustizia distributiva e commutativa: punire chi sbaglia e premiare il giusto. Ma non è solo questa la giustizia: ci sono anche le unità di misura e le regole per chi fa banca ed impresa (ultimamente a Siena se ne sono accorti…). La giustizia è fatta da queste due cose. Dalla giustizia nasce concordia, è la virtù di cui parla Lorenzetti: accorda gli animi. Ventiquattro cittadini si legano nella concordia alla giustizia. L’altro capo della corda che viene dalla giustizia ce l’ha in mano il popolo, e decide lui a chi darlo e a chi toglierlo. Lo da al comune, il bene comune: si vede il volto di Siena, la lupa con i gemelli e la scritta SCSV “Comune di Siena città della Vergine”. Il popolo gli da corda, e il popolo può toglierla. “E questi a cui un comun per loro si fanno lo qual, per governare qui stato, elegge di non tener giammai gli ochi rivolti da lo splendor de volti, de le virtù che ntorno a lui si stanno”: per governare il popolo il comune non deve mai togliere gli occhi dallo splendore delle virtù, fede, speranza e carità, quest’ultima virtù propria della politica. Oltre ad esse le quattro virtù cardinali, politiche per eccellenza. Più la magnanimità e la pace, aggiunte dall’autore. Pace è il centro dell affresco (la donna in bianco) perché è il cuore del desiderio dell’uomo. L’affresco è completato dai delinquenti, che non si legano al bene comune e quindi vanno legati e imprigionati.



Gli effetti del bene comune



Quando una società funziona così, gli si pagano volentieri anche le tasse. Tutti gli effetti buoni sono questi: l’affresco degli effetti buoni è sulla parete orientale del palazzo, dove sorge il sole: ci sono appunto i colori del sole che sorge. E’ una bellissima città, piena di gente, ma non una di quelle perfette città ideali del rinascimento, vuote. Si nota che è proprio Siena (nell’angolo in alto a sinistra compare il duomo). Il primo effetto è la bellezza della città. Il secondo effetto è che la città cresce (i muratori alzano case). il terzo effetto è che si lavora bene, e dove non si lavora si studia in università. E’ tutto un brulicare di lavoro, e avvicinandosi alle mura cresce lo scambio con la campagna. Bellezza, crescita, lavoro, studio, e, infine, matrimonio tra uomo e donna. La dove il mondo è governato dal bene comune nascono bambini. E in più ci si diverte, come si vede nelle feste e nei giochi. Anche Monte dei Paschi l’hanno costruita nel trecento: noi possiamo solo distruggere, come si è visto. Persino la campagna è bella e sicura: l’ultimo effetto di una politicante funziona è appunto la scurezza, la scurita. “Senza paura ognun franco cammini, e lavorando semini ciascuno”. La campagna è dominata dalla sicurezza. Vedendo questo mondo bellissimo si vede oltre la porta sottostante l’affresco (sempre aperta) un affresco raffigurante il paradiso, nella stanza accanto.



Il bene proprio



Nella parere occidentale è notte, tramonto. Anche il tempo ha voluto rovinare solo questo affresco. Qui c’è l’allegoria del bene proprio, del cattivo governo. Qui la giustizia è una donna piangente, spoglia e legata, tra i piatti spezzati della bilancia (in bianco). Nessuno si lega al bene comune e ognuno tira dalla sua parte, non da una parte comune, e perciò domina la tirannia. Il tiranno è una figura demoniaca, che ha come metodo la violenza, e la ricchezza come scopo. La tirannia però non è intesa come dittatura, ma è chi cerca il bene proprio! La caratteristica principale del tiranno è che è strabico, cieco. Il bene proprio è cieco perché nessuno può essere felice in un mondo infelice. Quindi chi fa il bene proprio è innanzitutto scemo. Ma quindi perché l’uomo cerca il bene proprio? Ci sono tre bestie: avarizia, superbia, vanagloria (ricerca del successo), e sei donne vestite di nero le accompagnano: crudeltà, inganno, tradimento, furore, divisione e guerra. L’ultimo particolare è che il tiranno non può tenere i piedi per terra, ma su quella brutta bestia che è l’amica dei tiranni: il caprone della lussuria (se avete pensato ad una persona sola vi sbagliate di grosso: sappiate che è ben accompagnato…).



Gli effetti del bene proprio

Gli effetti sulla città sono tutto il contrario di prima. La città è brutta decadente, le botteghe sono chiuse, tranne quella del fabbro, sempre in costruzione di armi e armature per la guerra. Anche il matrimonio è distrutto (che profezia anche questa…): una sposa viene rapita dal marito ucciso, che giace morto per terra tra l’indifferenza della gente. Uomini armati devastano la campagna, e su tutto vige il timore, la paura. Le nuvole nere sulla campagna tengono in mano l’ultimo cartiglio dell’autore: “Per volere il bene proprio la giustizia è sottomessa alla tirannia: per questa via c’è solo morte e ruberia”.




Ed oggi?

Ma non si può finire così. L’illusione è che la colpa sia tutta del politici: siamo sessanta milioni di persone per bene guidate da mille deficienti? Se il problema fossero solo quelli in parlamento sarebbe facilissimo: cambiarli. Magari, sarebbe facile. E invece non è così. “In un paese democratico il parlamento è fatto da 10% dei migliori uomini, il 10% dei peggiori uomini e il resto è come il paese”, disse Churchill. Da cosa è nato l’ideale di Siena? Mi sono imbattuta in una storia mentre scrivevo il mio libro. Seconda Guerra Mondiale, 1944: gli americani bombardano la periferia di Siena e alcune bombe colpiscono la basilica di San Bernardino. C’era un crocifisso del 1300 in quella chiesa: ora non esiste più, ma la cosa sorprendente fu che dei frati (con dei soldati brianzoli) si recarono il giorno dopo il bombardamento tra le macerie della chiesa cercando le ostie.  Trovarono la testa di quel crocifisso, che si aprì in due tra le mani di uno dei due frati. Dentro era cava, e conteneva un cartiglio! Per sette secoli fu nascosto lì, e solo una bomba americana lo fece arrivare a noi.  Anno, mese, autore:  gennaio 1337 (stesso anno e mese in cui Lorenzetti inizia i suoi lavori), Pietro da Siena. Questo cartiglio contiene una preghiera: questo è il segreto per cui nasce una città così: un uomo che nel sul lavoro quotidiano nasconde una preghiera a Dio per se, la sua famiglie, la sua generazione, e per tutti quanti. E’ questa la strada, non c’è scorciatoia, checché ne dica Grillo, che vuole solo distruggere il parlamento mandando a casa solo i politici.

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;
}

14/2/2013

Aggiornamento delle sezioni del sito. Ho aggiunto il link alla sezione Tutorial dove si possono trovare i tutorial di programmazione in C++, Html, Cuda e i tutorial di Grafica.

AIR
ARCHITECTURE
ARTISTIC
BOOK
CONCRETE
FABRIC
FLOOR
GROUND
LEATHER
MANMADE
METAL
NATURE
PAPER
PLASTIC
ROAD
ROCK
RUSTED
WOOD

Ho aggiunto le mie textures nella sezione "Textures" del sito, tutte in HD (2560x1920).
Sono divise per categorie e sono sempre in aggiornamento.
Scaricabili gratuitamente.

Le operazioni tra le matrici si può dire che siano il punto forte di CUDA: la scheda grafica infatti è appositamente costruita per gestire tutti i pixel dello schermo contemporaneamente, quindi una matrice ad esempio 1280x800. Si capisce che un calcolo effettuato in serie con una potente CPU può essere svolto molto più rapidamente in parallelo tramite una GPU, e anche se ogni processore della scheda grafica presenta una potenza di calcolo nettamente inferiore rispetto al processore principale del computer, il calcolo del prodotto degli elementi di due matrici punto per punto è solo una banale operazione di moltiplicazione, svolta quindi alla medesima velocità.
Cercheremo quindi in questo programma di far gestire alla scheda grafica tutta la matrice contemporaneamente (o almeno a blocchi), di modo da effettuare tutti i calcoli elemento per elemento in un solo singolo colpo di clock.


#include < iostream > 
#include < cstdlib >
#include < cmath >
#define EPS 1e-6
#include < ctime >
using namespace std;
#define MAX 30
#define THREADS 32
typedef float dato; // così­ modifico solo questo float/double se voglio cambiare scheda

// funzione per controllare la coerenza dei risultati cpu-gpu
void check(dato *a, dato *b,int N) {

for(int i = 0; i < N; ++i)
for(int j = 0 ; j < N ; ++j )
if(fabs(a[i*N+j]-b[i*N+j])>EPS) {
cerr << "Errore in posizione (" << i << "," << j <<"), a: "\
<< a[i*N+j] << ", check: " << b[i*N+j] << endl; 
return;
}
cerr << "OK" << endl;
return;
};

//Versione cpu del prodotto matrice matrice
void matrix_product_cpu(dato *a,dato *b, dato *c,int N){

dato temp;
for(int i = 0; i < N; ++i) 
for(int j = 0; j < N; ++j) {
temp = 0; 
for(int k = 0; k < N;++k) 
temp += a[i*N+k]*b[k*N+j];
c[i*N+j] = temp;
}
return;
};

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

// Prodotto matriciale: versione ottimizzata. 
// Si considerano matrici quadrate di ordine N  multiplo della dimensione del blocco di threads THREADS (quadrato).
// E' possibile ottimizzare il calcolo di un gruppo di righe per un gruppo di colonne
// specificamente di un blocco di elementi di A di dimensione THREADSxN e un blocco di B di dimensione Nx THREADS. 
// Per ridurre il numero di accessi alla memoria è conveniente sfruttare la memoria condivisa della scheda grafica.
// La memoria condivisa è presente in quantità limitata e quindi diventa necessario ridurre il problema in tanti sottoproblemi.
// Piu' precisamente, il sottoproblema elementare è costituito da un prodotto di due matrici di dimensione THREADSxTHREADS
// memorizzate sulla memoria veloce (shared) della GPU

__global__ void matrix_product( dato* A, dato* B, dato* C, int width) {

    // la mia posizione è data da due coordinate
    int ib = blockIdx.y;
    int jb = blockIdx.x;        
    int it = threadIdx.y;
    int jt = threadIdx.x;
        
    // Indice della prima sottomatrice di A elaborata dal blocco
    // width è un multiplo intero di THREADS (larghezza matrice completata)
    // aBegin  include un certo numero ib  di gruppi di blocchi rettangolari THREADSxwidth
    int aBegin = width*THREADS*ib; // posiz in alto a sx dell'ib-esimo blocco 
    
    // Indice dell'ultima sottomatrice di A elaborata dal blocco
    int aEnd   = aBegin + width - 1; // condizione che mi dica qual è l'ultimo elemento della riga per farci i conti
    
    // numero di colonne tra una sottomatrice e la successiva
    int aStep  = THREADS;
    
    // Indice della prima sottomatrice di B elaborata dal blocco
    // bBegin include un certo numero jb di blocchi di colonne, blocchi larghi THREADS 
    int bBegin = THREADS*jb;
         
    // numero di elementi tra una sottomatrice e la successiva    
    int bStep  = THREADS * width;
     
    // Csub è usata come variabile in cui memorizzare il valore dell'elemento di C
    // calcolato dal thread
    // Viene aggiornato ripetutamente nel ciclo for seguente
    dato Csub = 0;    
     
    // Iterazione sulle sottomatrici in cui viene suddiviso il calcolo degli elementi del blocco 
// a <= aEnd controllo solo sulle a, tanto è quadrata
// questo ciclo for mi fa saltare da una sottomatrice di a alla successiva, idem per b
    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {

        // Dichiarazione della variabile in cui salvare la sottomatrice di A in esame
__shared__ dato As[THREADS][THREADS];
        __shared__ dato Bs[THREADS][THREADS];

        // Carico gli elementi di ciascuna sottomatrice in memoria condivisa: ogni thread del 
// blocco carica un elemento!      
As[it][jt] = A[a +  width*it + jt]; 
Bs[it][jt] = B[b +  width*it + jt];

        // Sincronizzo per assicurarmi che ogni thread del blocco abbia caricato gli elementi.
__syncthreads();

  // Calcolo i contributi agli elementi di matrice di C dovute alle sottomatrici in esame        
for (int k = 0; k < THREADS ; ++k ) Csub += As[it][k]*Bs[k][jt];
// l'elemento C[it][jt] viene aggiornato in nblocks passaggi, pari al numero di iterazioni
// del ciclo for
  
        // Sincronizzo per assicurarmi che il calcolo precedente sia terminato prima di caricare nuove
// sottomatrici
__syncthreads();
    }

    // Scrivo i risultati in C. Ogni thread elabora un elemento di C. 
    int c = width*THREADS*ib + THREADS*jb; // conto quante righe ci sono prima + di quante mi sposto         
    C[c +  width*it + jt] = Csub;    
};

// costruisco una cornice di zeri per ignorare ciò che eccede la dimensione N*N della matrice originaria
__global__ void completa_matrice(dato *a, int N) {

int i = blockIdx.y*blockDim.y + threadIdx.y;
int j = blockIdx.x*blockDim.x + threadIdx.x;
int offset = gridDim.x*blockDim.x*i + j;
        
        // se sono fuori dalla matrice N*N pongo tutto a zero!
if(i >= N || j >= N) a[offset] = 0.;

};

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

void matrix(dato *a, dato *b, dato *c, int N) {

dato *a_dev,*b_dev,*c_dev;
int nblock = (N+THREADS-1)/THREADS; // se N = nblock è già un multiplo intero di threads questa operazione non fa nulla
int width = nblock*THREADS; // calcolo larghezza matrice completata

cout << "width=" << width << "  dim matrix=" << N << "\n";
        
        // alloco la memoria su cuda per le due matrici di input e per quella di output
cudaMalloc((void**)&a_dev, width*width*sizeof(dato));
cudaMalloc((void**)&b_dev, width*width*sizeof(dato));
cudaMalloc((void**)&c_dev, width*width*sizeof(dato));

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

// copio da host a device: copio sulla scheda le matrici create nel main
        // copia bidimensionale: voglio accederci con doppia indicizzazione
cudaMemcpy2D(a_dev, width*sizeof(dato), a, N*sizeof(dato), N*sizeof(dato), N, cudaMemcpyHostToDevice);
cudaMemcpy2D(b_dev, width*sizeof(dato), b, N*sizeof(dato), N*sizeof(dato), N, cudaMemcpyHostToDevice);
// a ha N righe di larghezza N*sizeof(dato)
// a_dev è la destinazione sulla scheda di W*W elementi (width*sizeof(dato))
             // la matrice di partenza ha N*N elementi, ma la matrice di arrivo ha W*W elementi: alloco W*W elementi

        // lancio il kernel cuda: blocchi lungo x e lungo y, e all´interno threads lungo x e lungo y
dim3 c_threads(THREADS,THREADS);
dim3 c_blocks(nblock,nblock);

dim3 threads(THREADS,THREADS);
dim3 blocks(nblock,nblock);

// ho messo in a_dev e in b_dev una cornicetta di zeri
if(N % THREADS != 0) {
completa_matrice<<>>(a_dev,N);
       completa_matrice<<>>(b_dev,N);
}
// blocks contiene nblock nblock, threads contiene THREADS THREADS  
matrix_product<<>>(a_dev,b_dev,c_dev,width);

// copio il risultato in c (sulla cpu)
cudaMemcpy2D(c, N*sizeof(dato), c_dev, width*sizeof(dato), N*sizeof(dato), N, cudaMemcpyDeviceToHost);
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
float elapsed;
cudaEventElapsedTime(&elapsed,start,stop);
cout << N << " " << elapsed/1000. << endl;

// dealloco
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaFree(a_dev);
cudaFree(b_dev);
cudaFree(c_dev);
};

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

int main(int argc, char **argv) {

cudaSetDevice(0);
const int N = atoi(argv[1]); // N è la dimensione della matrice
// metto le matrici in vettori di N*N elementi
dato *a_host = new dato[N*N];
dato *b_host = new dato[N*N];
dato *c_host = new dato[N*N];
dato *c_check = new dato[N*N];

srand48(time(NULL));

// riempio matrici a e b con i numeri random
for(int i = 0; i < N ;++i)
for(int j = 0; j < N;++j) {
int index = i*N+j;
a_host[index] = drand48()*MAX; // numeri tra 0 e 30 (define MAX)
b_host[index] = drand48()*MAX;
}

float elapsed_cpu;
matrix(a_host,b_host,c_host,N); // esegue la maggior parte delle operazioni, definita sopra

clock_t start,stop;
start = clock();
matrix_product_cpu(a_host,b_host,c_check,N);
stop = clock();
elapsed_cpu = (float)(stop - start)/CLOCKS_PER_SEC;
cout << elapsed_cpu << endl;

check(c_host,c_check,N);
delete[](a_host);
delete[](b_host);
delete[](c_host);
delete[](c_check);

return 0;
};

Per confermare la potenza del calcolo su scheda grafica mediante CUDA proponiamo un programma di calcolo del prodotto scalare di due vettori casuali di lunghezza data che lo esegue sia su GPU che su CPU. I risultati sono sorprendenti, soprattutto aumentando il numero di punti. Ovviamente una scheda grafica potente avrà prestazioni maggiori di un processore medio, ma un processore molto potente può eseguire calcoli più velocemente di una scheda grafica con poca memoria allocata. In ogni caso, per vettori molto lunghi la supremazia della scheda grafica risulta schiacciante.


// PRODOTTO SCALARE TRA DUE VETTORI OTTIMIZZATO
// implementare lettura vettori da file
// confronto tra CPU e GPU

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

#define imin(a,b) (a
const int N = 10000000;

const int threadsPerBlock = 1024; // grandezze per l'ottimizzazione del problema (alternativa al 128*128 di prima)
const int blocksPerGrid = imin(32, (N+threadsPerBlock-1)/threadsPerBlock); // trucco per ottimizzare il numero di blocchi
//const int blocksPerGrid = 1024;

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

__global__ void scalar_prod( float *a,  float *b,  float *c ){

  __shared__ float cache[threadsPerBlock];   // shared: interna e comune ai threads di un singolo blocco
  int tid = threadIdx.x + blockDim.x*blockIdx.x;
  int cacheIndex = threadIdx.x; // indice del thread
  float temp = 0;

  while(tid < N){
    temp += a[tid]*b[tid]; // incremento dell'indice della componente sommata (salto di blocco in blocco!)
        tid += blockDim.x*gridDim.x;
  };

  // il numero di elementi da valutare può essere maggiore del numero di thread * numero di blocchi
  // la somma (temp) corre sui diversi gruppi di blocchi combinando le componenti associate 
  // allo stesso valore di tid (modulo il numero totale di threads)

  // scopo del prodotto scalare è di restituire un numero; parte di questo numero viene calcolato sulla GPU
  // la somma parziale appena calcolata viene messa nella memoria shared del blocco in corrispondenza del valore
  // di threadIdx.x
  cache[cacheIndex] = temp; // 512 (threads) somme parziali per blocco, ogni blocco ha la sua cache[]
 // sono privati all'interno del blocco: non si mischiano tra di loro!
  __syncthreads(); // aspetto che tutti i thread del blocco abbiano finito

        // valuto la somma parziale dei risultati dei vari threads di un blocco
  int i=blockDim.x/2;
  while(i!=0){
    if (cacheIndex < i) cache[cacheIndex] += cache[cacheIndex+i];
        // questa somma combina i valori della metà superiore del blocco di threads con quelli della 
        // meta' inferiore e li associa a indici della meta' inferiore
    __syncthreads();  // attendo che tutti i threads abbiano effettuato la somma
    // tutti gli elementi rilevanti hanno indici da zero a meta' del blocco nella seconda iterazione si
    // sommano gli elementi del secondo quarto con quelli del primo quarto del blocco, e così via
    i/=2; // poichè i è una variabile intera, nel caso 1/2 = 0.5 viene restituito 0
  };

  if (cacheIndex ==0) c[blockIdx.x] = cache[0];
} // il risultato finale e' la somma parziale di un blocco

/////////////////////////////////////////////////////////////////////////////////////////////////////
// prodotto scalare su cpu (confronto dei tempi)
void sp_cpu_check(float *a, float *b, float &check){
int i;
  check=0;
  for (i=0; i
    check += a[i]*b[i];
  }
}

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

int main( void ){

  float *a, *b, *partial_c;
  float *dev_a, *dev_b, *dev_partial_c;
  float  sp_gpu = 0, sp_cpu = 0;

  a = (float *)malloc(N*sizeof(float));
  b = (float *)malloc(N*sizeof(float));
  partial_c = (float *)malloc(blocksPerGrid*sizeof(float));

  // riempimento casuale di vettori (escluso dal conteggio dei tempi)
  for( int i=0; i
    a[i] = drand48();
    b[i] = drand48();
  };

  // alloco la memoria sulla scheda
  cudaMalloc( (void **)&dev_a, N*sizeof(float) );
  cudaMalloc( (void **)&dev_b, N*sizeof(float) );
  cudaMalloc( (void **)&dev_partial_c, blocksPerGrid*sizeof(float) );

  // faccio partire il cronometro della gpu
  cudaEvent_t start1,stop1;
  float time1;
  cudaEventCreate(&start1);
  cudaEventCreate(&stop1);
  cudaEventRecord(start1,0);
  
  cudaMemcpy( dev_a, a, N*sizeof(float), cudaMemcpyHostToDevice );
  cudaMemcpy( dev_b, b, N*sizeof(float), cudaMemcpyHostToDevice );

  scalar_prod<<>>(dev_a, dev_b, dev_partial_c);

  cudaMemcpy( partial_c, dev_partial_c, blocksPerGrid*sizeof(float), cudaMemcpyDeviceToHost );

  // infine le somme parziali dei vari blocchi vengono combinate 
  for( int i=0; i
      sp_gpu += partial_c[i];
  };
cout << "scalar product gpu = " << sp_gpu << "\n";

  // fermo il cronometro della gpu
  cudaEventRecord(stop1,0);
  cudaEventSynchronize(stop1);
  cudaEventElapsedTime(&time1, start1, stop1);
  cout << "time gpu: " << time1 << "\n";
  
  // delloco dalla gpu
  cudaEventDestroy(start1);
  cudaEventDestroy(stop1);
  cudaFree(dev_a);
  cudaFree(dev_b);
  cudaFree(dev_partial_c);

  ///////////////////////////////////////////
  // rifaccio il calcolo con la cpu
  cudaEvent_t start2,stop2;
  float time2;
  cudaEventCreate(&start2);
  cudaEventCreate(&stop2);
  cudaEventRecord(start2,0);

  //check del prodotto scalare, calcolato sulla cpu
  sp_cpu_check(a,b,sp_cpu);
cout << "scalar product cpu = " << sp_cpu << "\n";

  // fermo il cronometro della cpu
  cudaEventRecord(stop2,0);
  cudaEventSynchronize(stop2);
  cudaEventElapsedTime(&time2, start2, stop2);
cout << "time cpu: " << time2 << "\n";

  // dealloco dalla cpu
  cudaEventDestroy(start2);
  cudaEventDestroy(stop2);
  free(a);
  free(b);
  free(partial_c);

  return 0;
}

Ora un programma ottimizzato per il calcolo del prodotto scalare tra vettori.


// PRODOTTO SCALARE TRA DUE VETTORI
// implementare lettura vettori da file

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

#define NBLOCKS 1024
#define THREADS 1024

const long int N = 100000;
//const int blocksPerGrid = imin( 32, (N+threadsPerBlock-1)/threadsPerBlock );

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

__global__ void scalar_prod( long int *a,  long int *b,  long int *c){

// la componente trattata viene individuata sommando un offset dovuto
  // al fatto di trovarsi in un certo blocco e quindi pari all'indice
    // del blocco per il numero di threads contenuti in un blocco
    // a cui si aggiunge l'indice del trhread corrente in quel blocco
    int tid = threadIdx.x + blockDim.x*blockIdx.x;

    // tid continua a venire incrementato, fino a quando non si eccede il numero massimo N di componenti

    while(tid < N){
      c[tid]=a[tid]*b[tid]; // salto un numero di threads pari a tutti quelli permessi lungo x    
      tid += blockDim.x*gridDim.x;
    }
    // questo while è un loop che mi permette di applicare il kernel ripetutamente (itero le 
    // 128*128 = 16.000 circa istanze = gridDim.x)
} 

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

int main( void ){

  long int a[N], b[N], c[N];
  long int *dev_a, *dev_b, *dev_c;
  long int  scalarprod = 0;

  // riempimento dei vettori secondo questa scelta
  for(int i=0; i
    a[i] = 2*i;
    b[i] = -i;
  };

  // alloco memoria sulla scheda
  cudaMalloc( (void **)&dev_a, N*sizeof(long int) );
  cudaMalloc( (void **)&dev_b, N*sizeof(long int) );
  cudaMalloc( (void **)&dev_c, N*sizeof(long int) );

  // faccio partire il cronometro (non conto il tempo di creazione dei vettori)
  cudaEvent_t start,stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);
  cudaEventRecord(start,0);

  cudaMemcpy( dev_a, a, N*sizeof(long int), cudaMemcpyHostToDevice );
  cudaMemcpy( dev_b, b, N*sizeof(long int), cudaMemcpyHostToDevice );

// scalar_prod<<<(N+127)/128,128>>>(dev_a, dev_b, dev_c);
   scalar_prod<<>>(dev_a, dev_b, dev_c);

  cudaMemcpy( c, dev_c, N*sizeof(long int), cudaMemcpyDeviceToHost ); // passo all'host
  // può ora essere grande come la memoria ram della scheda grafica!!

   for( int i=0; i
      cout << "scalar product = " << scalarprod << "\n";

  // fermo il cronometro (chiedo che la gpu abbia finito prima di restituire il val di stop
  cudaEventRecord(stop,0);
  cudaEventSynchronize(stop);
  float elapsed;
  cudaEventElapsedTime(&elapsed, start, stop);
  cout << "time: " << elapsed/1000. << endl;

  // dealloco
  cudaEventDestroy(start);
  cudaEventDestroy(stop);
  cudaFree(dev_a);
  cudaFree(dev_b);
  cudaFree(dev_c);

  return 0;
}

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;
}

Pratica: creazione di un programma in CUDA che esegue la somma di due numeri. Nei commenti la spiegazione del codice.
Ricordo che per utilizzare programmi in CUDA bisogna avere una scheda grafica nvidia CUDA capable e un compilatore.

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

// questa funzione mi dice di lavorare sulla scheda grafica
// ogni kernel viene dichiarato con __global__
__global__ void add( int a, int b, int *c){ 

  *c= a+b;
}

// nel main creo le variabili
int main( void ){

  int a,b,c;
  int *dev_c;
  a=3;
  b=5;

// 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 (-> non posso scrivere (int **) perché puó 
// essere che la cpu e la gpu leggano in modo diverso gli interi)
  cudaMalloc( (void **)&dev_c, sizeof(int)  );

// il kernel viene invocato e il risultato viene associato a dev_c
// 1,1 indica che sto usando un blocco solo con un thread solo 
// (altrimenti mi verrebbe 1024 volte il risultato!)
  add<<<1>>>(a, b, dev_c); 

// all'indirizzo di c viene copiato 
// quanto si trova sul device a partire da dev_c
// per una lunghezza pari a sizeof(int)

// dev_c non è accessibile dal main perchè va sulla gpu! 
// infatti devo copiarlo sulla cpu con questo codice
  cudaMemcpy( &c, dev_c, sizeof(int), cudaMemcpyDeviceToHost); //cudaMemcpyHostToDevice

  cout << "3+5=" << c << "\n";

// libero la memoria allocata da cudaMalloc
  cudaFree( dev_c);

  return 0;
}

La scheda grafica è un complesso componente hardware contenuto nella maggior parte dei dispositivi elettronici con una uscita video: nasce appunto con lo scopo di generare, a fronte di istruzioni della CPU, un segnale elettrico che possa essere inviato ad un display per poi essere tradotto in segnale ottico. La grande potenzialità delle schede grafiche, che le rende uno strumento molto utile anche per eseguire operazioni matematiche, risiede nella loro capacità intrinseca di eseguire un certo numero di operazioni contemporanea-
mente; devono infatti essere in grado di gestire contemporaneamente tutti i pixel dello schermo, modicandoli in tempo reale.
Tramite il linguaggio CUDA, sviluppato dalla NVIDIA sulle sue schede grafiche, è possibile scrivere programmi basati sul C++ che permettono di sfruttare la potenza del calcolo parallelo.

Il calcolo in parallelo puo rappresentare una soluzione intelligente per approcciare alcuni problemi computazionali, come ad esempio le simulazioni, dove un certo numero di operazioni uguali devono essere eseguite da un programma di calcolo.

Il linguaggio CUDA permette di velocizzare i processi parallelizzando le operazioni uguali, che vengono cos svolte in contemporanea dalla GPU, a fronte di istruzioni date dalla CPU.



Vediamo ora un primo esempio di programma scritto in CUDA:


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

__global__ void kernel( void ){ // questa funzione mi dice di lavorare sulla scheda grafica
}

int main(){
  cudaSetDevice(1); // scelgo la scheda grafica (in caso la macchina ne abbia più di una)
  kernel<<<1>>>();
  cout << "ciao\n";
  return 0;
}


Come possiamo vedere, le funzioni che devono girare sulla scheda grafica hanno bisogno dell'attributo __global__ prima della definizione, e al momento del lancio del kernel CUDA è necessario scegliere mediante il comando <<<1>>> il numero di threads e blocks su cui farla girare.

In questo caso il kernel CUDA è solo una funzione void, che però è lanciata non dalla CPU ma dalla GPU.


Per eseguire utilizziamo il seguente makefile:



LIBS=
#-lcuda -lm -L/usr/local/cuda/lib/ -lcudart #la uso solo se sono sulla macchina q
OBJS=
NVCC=nvcc

ECHO=/bin/echo

#default: hello.x
#default: devprop.x
#default: add.x
#default: somma_vec.x
#default: scalar_prod_blocchi.x
#default: scalar_prod_thread.x
default: scalar_prod_full.x

%.x: %.cu
$(NVCC) $(OBJS) $(LIBS) $< -o $@

clean:
rm -f *.x *.o

Ho aggiornato completamente il sito! Ora è suddiviso in varie sezioni, accessibili dalla prima pagina e sulla sidebar a destra. Illustro brevemente di che si tratta:

Home page: una slideshow dei miei scatti migliori come sfondo, due linee di link in alto e in basso per accedere alle sezioni più importanti. La home è raggiungibile tramite il logo in alto a sinistra.
Blog: questa è un po' la bacheca del sito, dove posterò le novità e le notizie in primo piano.
Foto: è l'archivio delle mie foto, caricate su siti di upload gratuiti. Presto saranno protette da password per la tutela della privacy, quindi se le desiderate chiedetemela!
Testi: appunti e testi, sbobinature e racconti. E' la sezione di cultura.
Tutorial: suddivisi in tutorial di C++, Cuda, Html, CSS e grafica, ognuno avrà la sua sezione e la sua etichetta.
Digital Art: il mio portfolio di lavori di grafica. Tags, pictures, disegni in computer graphic, disegni a mano libera, texture (divise per categoria), scatti migliori e video.

Nel footer, alcune informazioni personali e i contatti, oltre alla mappa del sito.

In questa sezione mostrerò alcuni tutorial sull'utilizzo di CUDA, estensione del C++ progettata da nvidia per eseguire programmi su scheda grafica.

Di seguito il link alla pagina ufficiale del progetto: link

In questa sezione propongo alcuni tutorial di web design di Html, CSS e qualcosina di javascript e php.

Le ragioni del voto

In base a che cosa devo andare a votare? Ne vale la pena? Che criteri e ragioni ho?

- Questo paese può avere un futuro se ce l'hanno i suoi giovani. Può sembrare ovvio, ma non fa notizia, purtroppo. Oggi il mondo ha aggiunto una direzione e una velocità per cui non si riesce a stare dietro, e sta accelerando. Un cambiamento e una accelerazione impone un ripensamento del nostro modo d'essere. Questo è un primo motivo che può portare i giovani a non considerare l'idea di mettersi in gioco: se i giovani non si riconoscono in ciò che succede al mondo non si riconoscono neanche nei loro diritti e doveri. La nostra società è in cambiamento: tra pochi anni sarà completamente diversa. Allora questo è il tempo della scelta: cosa fare, con chi stare. La x del voto non deve essere una x di non conoscenza, come le persone analfabete che si firmavano così: deve essere una x con coscienza. In un mondo complesso è difficile appartenere completamente a qualcosa: ad un partito, per esempio. Favola di esopo: non potremo mai essere categorizzati in un partito. In questa società così complessa la situazione è difficile: i leader non riescono a far convivere tutto, a far convivere le differenze e i paradossi. Facendo convivere le differenze si possono far partire le cose, e quindi il voto non è un annichilimento di noi stessi in un partito ma una scelta consapevole. Ma perché è imporante il voto e il parere dei giovani?

Perché il giovane ha la possibilità di essere intelligente, cioè di avere una visione aperta perché ha davanti la vita. I giovani hanno la grande responsabilità di alzarsi e di gridare terra anche se siamo in mezzo ad una tempesta. La forza del giovane è guardare attraverso, intellegere, e dare questo sguardo anche agli altri che giovani non sono. In università abbiamo aumentato le tasse per sistemare il bilancio, e non lo faremo più perché è a posto, ma lo abbiamo fatto con i giovani; i corsi li prepariamo con i giovani, perché hanno l'intelligenza di vedere oltre l'orizzonte. Un paese può essere forte se le istituzioni sono forti: la politica, la chiesa, l'università.

Fra: oggi come può un giovane fare politica?
- Fare politica significa non occuparsi solo di se stessi. Quando uno inizia a preoccuparsi anche degli altri sta facendo politica. È vero, poi c'è da fare il passo di candidarsi: ma se uno fa quel passo con questa consapevolezza è già eletto. Dobbiamo recuperare nella parola politica l'elemento altruistico e di gratuità. Chi fa capire che il voto è una scelta drammatica porta dentro di sé una visione distorta: chi vince non porta con sé il bottino della vittoria. Non è questa la cosa importante. Come la frase con cui Bush ha iniziato il discorso di saluto e di benvenuto ad Obama: "sale sui gradini del Campidoglio un uomo che rappresenta il futuro del nostro paese" È incredibile che due avversari politici si considerino così, nel mondo di oggi.

Cri: cosa c'entra il mio studio con le elezioni?
- Studiando un autore di 2000 anni fa significa confrontare quella storia con ciò che sta accadendo ora. Se non ci fosse stato quell'autore non avresti avuto quella possibilità. Io mi sono laureato in ingegneria nucleare quando l'Italia ha rinunciato al nucleare. Quindi ho iniziato a vivere reinventandomi, e i miei studi mi hanno portato verso questa consapevolezza: questa opportunità è significata il mio futuro. Dobbiamo costruire opportunità con le elezioni. Se vogliamo dare un furturo non dobbiamo dare un benessere, ma una opportunità: così la vita viene resa più bella.